[clang] 340fa3e - [OpenACC] Implement firstprivate lowering except init. (#153847)

via cfe-commits Mon, 18 Aug 2025 06:33:53 -0700

Author: Erich Keane
Date: 2025-08-18T06:33:40-07:00
New Revision: 340fa3e1bb723de53e9074f50aed13eb15820b47


URL: 
https://github.com/llvm/llvm-project/commit/340fa3e1bb723de53e9074f50aed13eb15820b47
DIFF: 
https://github.com/llvm/llvm-project/commit/340fa3e1bb723de53e9074f50aed13eb15820b47.diff

LOG: [OpenACC] Implement firstprivate lowering except init. (#153847)

This patch implements the basic lowering infrastructure, but does not
quite implement the copy initialization, which requires #153622.

It does however pass verification for the 'copy' section, which just
contains a yield.

Added: 
    clang/test/CIR/CodeGenOpenACC/combined-firstprivate-clause.cpp
    clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause-templates.cpp
    clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.cpp

Modified: 
    clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
    clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp

Removed: 
    


################################################################################
diff  --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp 
b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index 9194b522114bc..72e2c533254c9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -387,6 +387,20 @@ class OpenACCClauseCIREmitter final
     return recipeName;
   }
 
+  void createFirstprivateRecipeCopy(
+      mlir::Location loc, mlir::Location locEnd, mlir::Value mainOp,
+      CIRGenFunction::AutoVarEmission tempDeclEmission,
+      mlir::acc::FirstprivateRecipeOp recipe, const VarDecl *varRecipe,
+      const VarDecl *temporary) {
+    builder.createBlock(&recipe.getCopyRegion(), recipe.getCopyRegion().end(),
+                        {mainOp.getType(), mainOp.getType()}, {loc, loc});
+    builder.setInsertionPointToEnd(&recipe.getCopyRegion().back());
+
+    // TODO: OpenACC: Implement this copy to actually do something.
+
+    mlir::acc::YieldOp::create(builder, locEnd);
+  }
+
   // Create the 'init' section of the recipe, including the 'copy' section for
   // 'firstprivate'.
   template <typename RecipeTy>
@@ -401,12 +415,6 @@ class OpenACCClauseCIREmitter final
       cgf.cgm.errorNYI(exprRange, "OpenACC Reduction recipe init");
     }
 
-    if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp>) {
-      // We haven't implemented the 'init'/copy recipe for firstprivate yet, so
-      // NYI it.
-      cgf.cgm.errorNYI(exprRange, "OpenACC firstprivate recipe init");
-    }
-
     CIRGenFunction::AutoVarEmission tempDeclEmission{
         CIRGenFunction::AutoVarEmission::invalid()};
 
@@ -442,17 +450,12 @@ class OpenACCClauseCIREmitter final
     mlir::acc::YieldOp::create(builder, locEnd);
 
     if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp>) {
-      if (!varRecipe->getInit()) {
-        // If we don't have any initialization recipe, we failed during Sema to
-        // initialize this correctly. If we disable the
-        // Sema::TentativeAnalysisScopes in SemaOpenACC::CreateInitRecipe, 
it'll
-        // emit an error to tell us.  However, emitting those errors during
-        // production is a violation of the standard, so we cannot do them.
-        cgf.cgm.errorNYI(
-            exprRange, "firstprivate copy-init recipe not properly generated");
-      }
-
-      cgf.cgm.errorNYI(exprRange, "firstprivate copy section generation");
+      // TODO: OpenACC: we should have a errorNYI call here if
+      // !varRecipe->getInit(), but as that generation isn't currently
+      // implemented, it ends up being too noisy. So when we implement 
copy-init
+      // generation both in Sema and here, we should have a diagnostic here.
+      createFirstprivateRecipeCopy(loc, locEnd, mainOp, tempDeclEmission,
+                                   recipe, varRecipe, temporary);
     }
 
     // Make sure we cleanup after ourselves here.
@@ -1155,6 +1158,43 @@ class OpenACCClauseCIREmitter final
       llvm_unreachable("Unknown construct kind in VisitPrivateClause");
     }
   }
+
+  void VisitFirstPrivateClause(const OpenACCFirstPrivateClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp,
+                               mlir::acc::SerialOp>) {
+      for (const auto [varExpr, varRecipe] :
+           llvm::zip_equal(clause.getVarList(), clause.getInitRecipes())) {
+        CIRGenFunction::OpenACCDataOperandInfo opInfo =
+            cgf.getOpenACCDataOperandInfo(varExpr);
+        auto firstPrivateOp = mlir::acc::FirstprivateOp::create(
+            builder, opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
+            /*implicit=*/false, opInfo.name, opInfo.bounds);
+
+        firstPrivateOp.setDataClause(mlir::acc::DataClause::acc_firstprivate);
+
+        {
+          mlir::OpBuilder::InsertionGuard guardCase(builder);
+          auto recipe = getOrCreateRecipe<mlir::acc::FirstprivateRecipeOp>(
+              cgf.getContext(), varExpr, varRecipe.RecipeDecl,
+              varRecipe.InitFromTemporary,
+              Decl::castToDeclContext(cgf.curFuncDecl), opInfo.baseType,
+              firstPrivateOp.getResult());
+
+          // TODO: OpenACC: The dialect is going to change in the near future 
to
+          // have these be on a 
diff erent operation, so when that changes, we
+          // probably need to change these here.
+          operation.addFirstPrivatization(builder.getContext(), firstPrivateOp,
+                                          recipe);
+        }
+      }
+    } else if constexpr (isCombinedType<OpTy>) {
+      // Unlike 'private', 'firstprivate' applies to the compute op, not the
+      // loop op.
+      applyToComputeOp(clause);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitFirstPrivateClause");
+    }
+  }
 };
 
 template <typename OpTy>

diff  --git a/clang/test/CIR/CodeGenOpenACC/combined-firstprivate-clause.cpp 
b/clang/test/CIR/CodeGenOpenACC/combined-firstprivate-clause.cpp
new file mode 100644
index 0000000000000..6d15abc2fefd4
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/combined-firstprivate-clause.cpp
@@ -0,0 +1,571 @@
+// RUN: %clang_cc1 -fopenacc -triple x86_64-linux-gnu 
-Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple 
x86_64-linux-pc %s -o - | FileCheck %s
+
+struct NoCopyConstruct {};
+
+struct CopyConstruct {
+  CopyConstruct() = default;
+  CopyConstruct(const CopyConstruct&);
+};
+
+struct NonDefaultCtor {
+  NonDefaultCtor();
+};
+
+struct HasDtor {
+  ~HasDtor();
+};
+
+// CHECK: acc.firstprivate.recipe @firstprivatization__ZTSA5_7HasDtor : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> 
{{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!rec_HasDtor x 5>, 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> 
{{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+//
+// CHECK-NEXT: } destroy {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> 
{{.*}}):
+// CHECK-NEXT: %[[LAST_IDX:.*]] = cir.const #cir.int<4> : !u64i
+// CHECK-NEXT: %[[ARRPTR:.*]] = cir.cast(array_to_ptrdecay, %[[ARG]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>), !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: %[[ELEM:.*]] = cir.ptr_stride(%[[ARRPTR]] : 
!cir.ptr<!rec_HasDtor>, %[[LAST_IDX]] : !u64i), !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !cir.ptr<!rec_HasDtor>, 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, ["__array_idx"]
+// CHECK-NEXT: cir.store %[[ELEM]], %[[ITR]] : !cir.ptr<!rec_HasDtor>, 
!cir.ptr<!cir.ptr<!rec_HasDtor>>
+// CHECK-NEXT: cir.do {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ELEM_LOAD]]) nothrow : 
(!cir.ptr<!rec_HasDtor>) -> ()
+// CHECK-NEXT: %[[NEG_ONE:.*]] =  cir.const #cir.int<-1> : !s64i
+// CHECK-NEXT: %[[PREVELEM:.*]] = cir.ptr_stride(%[[ELEM_LOAD]] : 
!cir.ptr<!rec_HasDtor>, %[[NEG_ONE]] : !s64i), !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: cir.store %[[PREVELEM]], %[[ITR]] : !cir.ptr<!rec_HasDtor>, 
!cir.ptr<!cir.ptr<!rec_HasDtor>>
+// CHECK-NEXT: cir.yield
+// CHECK-NEXT: } while {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(ne, %[[ELEM_LOAD]], %[[ARRPTR]]) : 
!cir.ptr<!rec_HasDtor>, !cir.bool
+// CHECK-NEXT: cir.condition(%[[CMP]])
+// CHECK-NEXT: }
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTSA5_14NonDefaultCtor : 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> 
{{.*}}):
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 
5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor 
x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> 
{{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTSA5_13CopyConstruct : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> 
{{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!rec_CopyConstruct x 5>, 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 
5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> 
{{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTSA5_15NoCopyConstruct : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 
5>> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct 
x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> 
{{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSA5_f : 
!cir.ptr<!cir.array<!cir.float x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!cir.float x 5>, 
!cir.ptr<!cir.array<!cir.float x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> 
{{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSA5_i : 
!cir.ptr<!cir.array<!s32i x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 
5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS7HasDtor : 
!cir.ptr<!rec_HasDtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } destroy {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ARG]]) nothrow : 
(!cir.ptr<!rec_HasDtor>) -> ()
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, 
!cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS13CopyConstruct 
: !cir.ptr<!rec_CopyConstruct> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTS15NoCopyConstruct : !cir.ptr<!rec_NoCopyConstruct> init 
{
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_NoCopyConstruct, 
!cir.ptr<!rec_NoCopyConstruct>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSf : 
!cir.ptr<!cir.float> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.float> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.float, !cir.ptr<!cir.float>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.float> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!cir.float> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSi : 
!cir.ptr<!s32i> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):
+// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!s32i> {{.*}}, %[[ARG_TO:.*]]: 
!cir.ptr<!s32i> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+
+extern "C" void acc_combined() {
+  // CHECK: cir.func{{.*}} @acc_combined() {
+
+  int someInt;
+  // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, 
["someInt"]
+  float someFloat;
+  // CHECK-NEXT: %[[SOMEFLOAT:.*]] = cir.alloca !cir.float, 
!cir.ptr<!cir.float>, ["someFloat"]
+  NoCopyConstruct noCopy;
+  // CHECK-NEXT: %[[NOCOPY:.*]] = cir.alloca !rec_NoCopyConstruct, 
!cir.ptr<!rec_NoCopyConstruct>, ["noCopy"]
+  CopyConstruct hasCopy;
+  // CHECK-NEXT: %[[HASCOPY:.*]] = cir.alloca !rec_CopyConstruct, 
!cir.ptr<!rec_CopyConstruct>, ["hasCopy"]
+  NonDefaultCtor notDefCtor;
+  // CHECK-NEXT: %[[NOTDEFCTOR:.*]] = cir.alloca !rec_NonDefaultCtor, 
!cir.ptr<!rec_NonDefaultCtor>, ["notDefCtor", init]
+  HasDtor dtor;
+  // CHECK-NEXT: %[[DTOR:.*]] = cir.alloca !rec_HasDtor, 
!cir.ptr<!rec_HasDtor>, ["dtor"]
+  int someIntArr[5];
+  // CHECK-NEXT: %[[INTARR:.*]] = cir.alloca !cir.array<!s32i x 5>, 
!cir.ptr<!cir.array<!s32i x 5>>, ["someIntArr"]
+  float someFloatArr[5];
+  // CHECK-NEXT: %[[FLOATARR:.*]] = cir.alloca !cir.array<!cir.float x 5>, 
!cir.ptr<!cir.array<!cir.float x 5>>, ["someFloatArr"]
+  NoCopyConstruct noCopyArr[5];
+  // CHECK-NEXT: %[[NOCOPYARR:.*]] = cir.alloca 
!cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct 
x 5>>, ["noCopyArr"]
+  CopyConstruct hasCopyArr[5];
+  // CHECK-NEXT: %[[HASCOPYARR:.*]] = cir.alloca !cir.array<!rec_CopyConstruct 
x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["hasCopyArr"]
+  NonDefaultCtor notDefCtorArr[5];
+  // CHECK-NEXT: %[[NOTDEFCTORARR:.*]] = cir.alloca 
!cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 
5>>, ["notDefCtorArr", init]
+  HasDtor dtorArr[5];
+  // CHECK-NEXT: %[[DTORARR:.*]] = cir.alloca !cir.array<!rec_HasDtor x 5>, 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["dtorArr"]
+  // CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[NOTDEFCTOR]]) : 
(!cir.ptr<!rec_NonDefaultCtor>) -> ()
+
+#pragma acc parallel loop firstprivate(someInt)
+  for(int i = 0; i < 5; ++i);
+  // CHECK: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[SOMEINT]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSi -> %[[PRIVATE]] : !cir.ptr<!s32i>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(someFloat)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[SOMEFLOAT]] : 
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTSf -> %[[PRIVATE]] : !cir.ptr<!cir.float>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel loop firstprivate(noCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPY]] : 
!cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = 
"noCopy"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTS15NoCopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!rec_NoCopyConstruct>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(hasCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPY]] : 
!cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTS13CopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!rec_CopyConstruct>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(notDefCtor)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTOR]] : 
!cir.ptr<!rec_NonDefaultCtor>) -> !cir.ptr<!rec_NonDefaultCtor> {name = 
"notDefCtor"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTS14NonDefaultCtor -> %[[PRIVATE]] : 
!cir.ptr<!rec_NonDefaultCtor>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(dtor)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTOR]] : 
!cir.ptr<!rec_HasDtor>) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTS7HasDtor -> %[[PRIVATE]] : 
!cir.ptr<!rec_HasDtor>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel loop firstprivate(someInt, someFloat, noCopy, hasCopy, 
notDefCtor, dtor)
+  for(int i = 0; i < 5; ++i);
+  // CHECK: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[SOMEINT]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[SOMEFLOAT]] : 
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPY]] : 
!cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = 
"noCopy"}
+  // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPY]] : 
!cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTOR]] : 
!cir.ptr<!rec_NonDefaultCtor>) -> !cir.ptr<!rec_NonDefaultCtor> {name = 
"notDefCtor"}
+  // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTOR]] : 
!cir.ptr<!rec_HasDtor>) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}
+  // CHECK: acc.parallel combined(loop) firstprivate(@firstprivatization__ZTSi 
-> %[[PRIVATE1]] : !cir.ptr<!s32i>,
+  // CHECK-SAME: @firstprivatization__ZTSf -> %[[PRIVATE2]] : 
!cir.ptr<!cir.float>,
+  // CHECK-SAME: @firstprivatization__ZTS15NoCopyConstruct -> %[[PRIVATE3]] : 
!cir.ptr<!rec_NoCopyConstruct>,
+  // CHECK-SAME: @firstprivatization__ZTS13CopyConstruct -> %[[PRIVATE4]] : 
!cir.ptr<!rec_CopyConstruct>,
+  // CHECK-SAME: @firstprivatization__ZTS14NonDefaultCtor -> %[[PRIVATE5]] : 
!cir.ptr<!rec_NonDefaultCtor>,
+  // CHECK-SAME: @firstprivatization__ZTS7HasDtor -> %[[PRIVATE6]] : 
!cir.ptr<!rec_HasDtor>)
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc serial loop firstprivate(someIntArr[1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTSA5_i -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!s32i x 5>>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(someFloatArr[1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_f -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(noCopyArr[1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(hasCopyArr[1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(notDefCtorArr[1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(dtorArr[1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_7HasDtor -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(someIntArr[1], someFloatArr[1], 
noCopyArr[1], hasCopyArr[1], notDefCtorArr[1], dtorArr[1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTSA5_i -> %[[PRIVATE1]] : 
!cir.ptr<!cir.array<!s32i x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_f -> %[[PRIVATE2]] : 
!cir.ptr<!cir.array<!cir.float x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE3]] 
: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE4]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE5]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_7HasDtor -> %[[PRIVATE6]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>)
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel loop firstprivate(someIntArr[1:1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_i -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!s32i x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(someFloatArr[1:1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTSA5_f -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(noCopyArr[1:1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop firstprivate(hasCopyArr[1:1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}
+  // CHECK-NEXT: acc.serial combined(loop) 
firstprivate(@firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) {
+  // CHECK-NEXT: acc.loop combined(serial)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(notDefCtorArr[1:1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(dtorArr[1:1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_7HasDtor -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) {
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel loop firstprivate(someIntArr[1:1], someFloatArr[1:1], 
noCopyArr[1:1], hasCopyArr[1:1], notDefCtorArr[1:1], dtorArr[1:1])
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}
+  // CHECK-NEXT: acc.parallel combined(loop) 
firstprivate(@firstprivatization__ZTSA5_i -> %[[PRIVATE1]] : 
!cir.ptr<!cir.array<!s32i x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_f -> %[[PRIVATE2]] : 
!cir.ptr<!cir.array<!cir.float x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE3]] 
: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE4]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE5]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_7HasDtor -> %[[PRIVATE6]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>)
+  // CHECK-NEXT: acc.loop combined(parallel)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+}

diff  --git 
a/clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause-templates.cpp 
b/clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause-templates.cpp
new file mode 100644
index 0000000000000..a9f0dd99e3bd4
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause-templates.cpp
@@ -0,0 +1,90 @@
+// RUN: %clang_cc1 -fopenacc -triple x86_64-linux-gnu 
-Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple 
x86_64-linux-pc %s -o - | FileCheck %s
+
+struct CopyConstruct {
+  CopyConstruct() = default;
+  CopyConstruct(const CopyConstruct&);
+};
+
+struct NonDefaultCtor {
+  NonDefaultCtor();
+};
+
+struct HasDtor {
+  ~HasDtor();
+};
+
+// CHECK: acc.firstprivate.recipe @firstprivatization__ZTSi : !cir.ptr<!s32i> 
init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):
+// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!s32i> {{.*}}, %[[ARG_TO:.*]]: 
!cir.ptr<!s32i> {{.*}}):
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS7HasDtor : 
!cir.ptr<!rec_HasDtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } destroy {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ARG]]) nothrow : 
(!cir.ptr<!rec_HasDtor>) -> ()
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, 
!cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS13CopyConstruct 
: !cir.ptr<!rec_CopyConstruct> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+
+template<typename T, typename U, typename V, typename W>
+void dependent_version(const T &cc, const U &ndc, const V &dtor, const W 
&someInt) {
+  // CHECK: cir.func 
{{.*}}@_Z17dependent_versionI13CopyConstruct14NonDefaultCtor7HasDtoriEvRKT_RKT0_RKT1_RKT2_(%[[ARG0:.*]]:
 !cir.ptr<!rec_CopyConstruct> {{.*}}, %[[ARG1:.*]]: 
!cir.ptr<!rec_NonDefaultCtor> {{.*}}, %[[ARG2:.*]]: !cir.ptr<!rec_HasDtor> 
{{.*}}, %[[ARG3:.*]]: !cir.ptr<!s32i> {{.*}}) {
+  // CHECK-NEXT: %[[CC:.*]] = cir.alloca !cir.ptr<!rec_CopyConstruct>, 
!cir.ptr<!cir.ptr<!rec_CopyConstruct>>, ["cc", init, const]
+  // CHECK-NEXT: %[[NDC:.*]] = cir.alloca !cir.ptr<!rec_NonDefaultCtor>, 
!cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, ["ndc", init, const]
+  // CHECK-NEXT: %[[DTOR:.*]] = cir.alloca !cir.ptr<!rec_HasDtor>, 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, ["dtor", init, const]
+  // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !cir.ptr<!s32i>, 
!cir.ptr<!cir.ptr<!s32i>>, ["someInt", init, const]
+  //             %             3 = cir.alloca !cir.ptr<!s32i>, 
!cir.ptr<!cir.ptr<!s32i>>, ["someInt", init, const]
+
+#pragma acc parallel firstprivate(cc, ndc, dtor, someInt)
+  ;
+  // CHECK: %[[PRIV_LOAD:.*]] = cir.load %[[CC]] : 
!cir.ptr<!cir.ptr<!rec_CopyConstruct>>, !cir.ptr<!rec_CopyConstruct>
+  // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[PRIV_LOAD]] : 
!cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "cc"}
+  // CHECK-NEXT: %[[PRIV_LOAD:.*]] = cir.load %[[NDC]] : 
!cir.ptr<!cir.ptr<!rec_NonDefaultCtor>>, !cir.ptr<!rec_NonDefaultCtor>
+  // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[PRIV_LOAD]] : 
!cir.ptr<!rec_NonDefaultCtor>) -> !cir.ptr<!rec_NonDefaultCtor> {name = "ndc"}
+  // CHECK-NEXT: %[[PRIV_LOAD:.*]] = cir.load %[[DTOR]] : 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, !cir.ptr<!rec_HasDtor>
+  // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[PRIV_LOAD]] : 
!cir.ptr<!rec_HasDtor>) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}
+  // CHECK-NEXT: %[[PRIV_LOAD:.*]] = cir.load %[[SOMEINT]] : 
!cir.ptr<!cir.ptr<!s32i>>, !cir.ptr<!s32i>
+  // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[PRIV_LOAD]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+
+  // CHECK-NEXT: acc.parallel 
firstprivate(@firstprivatization__ZTS13CopyConstruct -> %[[PRIVATE1]] : 
!cir.ptr<!rec_CopyConstruct>, 
+  // CHECK-SAME: @firstprivatization__ZTS14NonDefaultCtor -> %[[PRIVATE2]] : 
!cir.ptr<!rec_NonDefaultCtor>,
+  // CHECK-SAME: @firstprivatization__ZTS7HasDtor -> %[[PRIVATE3]] : 
!cir.ptr<!rec_HasDtor>,
+  // CHECK-SAME: @firstprivatization__ZTSi -> %[[PRIVATE4]] : !cir.ptr<!s32i>) 
{
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+}
+
+void use() {
+  CopyConstruct cc;
+  NonDefaultCtor ndc;
+  HasDtor dtor;
+  int i;
+  dependent_version(cc, ndc, dtor, i);
+}

diff  --git a/clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.cpp 
b/clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.cpp
new file mode 100644
index 0000000000000..d25208c65ac20
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.cpp
@@ -0,0 +1,508 @@
+// RUN: %clang_cc1 -fopenacc -triple x86_64-linux-gnu 
-Wno-openacc-self-if-potential-conflict -emit-cir -fclangir -triple 
x86_64-linux-pc %s -o - | FileCheck %s
+
+struct NoCopyConstruct {};
+
+struct CopyConstruct {
+  CopyConstruct() = default;
+  CopyConstruct(const CopyConstruct&);
+};
+
+struct NonDefaultCtor {
+  NonDefaultCtor();
+};
+
+struct HasDtor {
+  ~HasDtor();
+};
+
+// CHECK: acc.firstprivate.recipe @firstprivatization__ZTSA5_7HasDtor : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> 
{{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!rec_HasDtor x 5>, 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> 
{{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+//
+// CHECK-NEXT: } destroy {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_HasDtor x 5>> 
{{.*}}):
+// CHECK-NEXT: %[[LAST_IDX:.*]] = cir.const #cir.int<4> : !u64i
+// CHECK-NEXT: %[[ARRPTR:.*]] = cir.cast(array_to_ptrdecay, %[[ARG]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>), !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: %[[ELEM:.*]] = cir.ptr_stride(%[[ARRPTR]] : 
!cir.ptr<!rec_HasDtor>, %[[LAST_IDX]] : !u64i), !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: %[[ITR:.*]] = cir.alloca !cir.ptr<!rec_HasDtor>, 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, ["__array_idx"]
+// CHECK-NEXT: cir.store %[[ELEM]], %[[ITR]] : !cir.ptr<!rec_HasDtor>, 
!cir.ptr<!cir.ptr<!rec_HasDtor>>
+// CHECK-NEXT: cir.do {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ELEM_LOAD]]) nothrow : 
(!cir.ptr<!rec_HasDtor>) -> ()
+// CHECK-NEXT: %[[NEG_ONE:.*]] =  cir.const #cir.int<-1> : !s64i
+// CHECK-NEXT: %[[PREVELEM:.*]] = cir.ptr_stride(%[[ELEM_LOAD]] : 
!cir.ptr<!rec_HasDtor>, %[[NEG_ONE]] : !s64i), !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: cir.store %[[PREVELEM]], %[[ITR]] : !cir.ptr<!rec_HasDtor>, 
!cir.ptr<!cir.ptr<!rec_HasDtor>>
+// CHECK-NEXT: cir.yield
+// CHECK-NEXT: } while {
+// CHECK-NEXT: %[[ELEM_LOAD:.*]] = cir.load %[[ITR]] : 
!cir.ptr<!cir.ptr<!rec_HasDtor>>, !cir.ptr<!rec_HasDtor>
+// CHECK-NEXT: %[[CMP:.*]] = cir.cmp(ne, %[[ELEM_LOAD]], %[[ARRPTR]]) : 
!cir.ptr<!rec_HasDtor>, !cir.bool
+// CHECK-NEXT: cir.condition(%[[CMP]])
+// CHECK-NEXT: }
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTSA5_14NonDefaultCtor : 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> 
{{.*}}):
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !cir.array<!rec_NonDefaultCtor x 
5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor 
x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> 
{{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTSA5_13CopyConstruct : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> 
{{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!rec_CopyConstruct x 5>, 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 
5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_CopyConstruct x 5>> 
{{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTSA5_15NoCopyConstruct : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 
5>> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!rec_NoCopyConstruct x 5>, 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct 
x 5>> {{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> 
{{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSA5_f : 
!cir.ptr<!cir.array<!cir.float x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!cir.float x 5>, 
!cir.ptr<!cir.array<!cir.float x 5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> 
{{.*}}, %[[ARG_TO:.*]]: !cir.ptr<!cir.array<!cir.float x 5>> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSA5_i : 
!cir.ptr<!cir.array<!s32i x 5>> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.array<!s32i x 5>, !cir.ptr<!cir.array<!s32i x 
5>>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FORM:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!cir.array<!s32i x 5>> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS7HasDtor : 
!cir.ptr<!rec_HasDtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_HasDtor, !cir.ptr<!rec_HasDtor>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } destroy {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_HasDtor> {{.*}}):
+// CHECK-NEXT: cir.call @_ZN7HasDtorD1Ev(%[[ARG]]) nothrow : 
(!cir.ptr<!rec_HasDtor>) -> ()
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: %[[ALLOCA:.*]] = cir.alloca !rec_NonDefaultCtor, 
!cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTS13CopyConstruct 
: !cir.ptr<!rec_CopyConstruct> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_CopyConstruct, !cir.ptr<!rec_CopyConstruct>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_CopyConstruct> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe 
@firstprivatization__ZTS15NoCopyConstruct : !cir.ptr<!rec_NoCopyConstruct> init 
{
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_NoCopyConstruct, 
!cir.ptr<!rec_NoCopyConstruct>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!rec_NoCopyConstruct> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSf : 
!cir.ptr<!cir.float> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!cir.float> {{.*}}):
+// CHECK-NEXT: cir.alloca !cir.float, !cir.ptr<!cir.float>, 
["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!cir.float> {{.*}}, 
%[[ARG_TO:.*]]: !cir.ptr<!cir.float> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.firstprivate.recipe @firstprivatization__ZTSi : 
!cir.ptr<!s32i> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!s32i> {{.*}}):
+// CHECK-NEXT: cir.alloca !s32i, !cir.ptr<!s32i>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: } copy {
+// CHECK-NEXT: ^bb0(%[[ARG_FROM:.*]]: !cir.ptr<!s32i> {{.*}}, %[[ARG_TO:.*]]: 
!cir.ptr<!s32i> {{.*}}):
+//
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+
+extern "C" void acc_compute() {
+  // CHECK: cir.func{{.*}} @acc_compute() {
+
+  int someInt;
+  // CHECK-NEXT: %[[SOMEINT:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, 
["someInt"]
+  float someFloat;
+  // CHECK-NEXT: %[[SOMEFLOAT:.*]] = cir.alloca !cir.float, 
!cir.ptr<!cir.float>, ["someFloat"]
+  NoCopyConstruct noCopy;
+  // CHECK-NEXT: %[[NOCOPY:.*]] = cir.alloca !rec_NoCopyConstruct, 
!cir.ptr<!rec_NoCopyConstruct>, ["noCopy"]
+  CopyConstruct hasCopy;
+  // CHECK-NEXT: %[[HASCOPY:.*]] = cir.alloca !rec_CopyConstruct, 
!cir.ptr<!rec_CopyConstruct>, ["hasCopy"]
+  NonDefaultCtor notDefCtor;
+  // CHECK-NEXT: %[[NOTDEFCTOR:.*]] = cir.alloca !rec_NonDefaultCtor, 
!cir.ptr<!rec_NonDefaultCtor>, ["notDefCtor", init]
+  HasDtor dtor;
+  // CHECK-NEXT: %[[DTOR:.*]] = cir.alloca !rec_HasDtor, 
!cir.ptr<!rec_HasDtor>, ["dtor"]
+  int someIntArr[5];
+  // CHECK-NEXT: %[[INTARR:.*]] = cir.alloca !cir.array<!s32i x 5>, 
!cir.ptr<!cir.array<!s32i x 5>>, ["someIntArr"]
+  float someFloatArr[5];
+  // CHECK-NEXT: %[[FLOATARR:.*]] = cir.alloca !cir.array<!cir.float x 5>, 
!cir.ptr<!cir.array<!cir.float x 5>>, ["someFloatArr"]
+  NoCopyConstruct noCopyArr[5];
+  // CHECK-NEXT: %[[NOCOPYARR:.*]] = cir.alloca 
!cir.array<!rec_NoCopyConstruct x 5>, !cir.ptr<!cir.array<!rec_NoCopyConstruct 
x 5>>, ["noCopyArr"]
+  CopyConstruct hasCopyArr[5];
+  // CHECK-NEXT: %[[HASCOPYARR:.*]] = cir.alloca !cir.array<!rec_CopyConstruct 
x 5>, !cir.ptr<!cir.array<!rec_CopyConstruct x 5>>, ["hasCopyArr"]
+  NonDefaultCtor notDefCtorArr[5];
+  // CHECK-NEXT: %[[NOTDEFCTORARR:.*]] = cir.alloca 
!cir.array<!rec_NonDefaultCtor x 5>, !cir.ptr<!cir.array<!rec_NonDefaultCtor x 
5>>, ["notDefCtorArr", init]
+  HasDtor dtorArr[5];
+  // CHECK-NEXT: %[[DTORARR:.*]] = cir.alloca !cir.array<!rec_HasDtor x 5>, 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>, ["dtorArr"]
+  // CHECK-NEXT: cir.call @_ZN14NonDefaultCtorC1Ev(%[[NOTDEFCTOR]]) : 
(!cir.ptr<!rec_NonDefaultCtor>) -> ()
+
+#pragma acc parallel firstprivate(someInt)
+  ;
+  // CHECK: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[SOMEINT]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: acc.parallel firstprivate(@firstprivatization__ZTSi -> 
%[[PRIVATE]] : !cir.ptr<!s32i>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(someFloat)
+  ;
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[SOMEFLOAT]] : 
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: acc.serial firstprivate(@firstprivatization__ZTSf -> 
%[[PRIVATE]] : !cir.ptr<!cir.float>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel firstprivate(noCopy)
+  ;
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPY]] : 
!cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = 
"noCopy"}
+  // CHECK-NEXT: acc.parallel 
firstprivate(@firstprivatization__ZTS15NoCopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!rec_NoCopyConstruct>
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(hasCopy)
+  ;
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPY]] : 
!cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: acc.serial 
firstprivate(@firstprivatization__ZTS13CopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!rec_CopyConstruct>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(notDefCtor)
+  ;
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTOR]] : 
!cir.ptr<!rec_NonDefaultCtor>) -> !cir.ptr<!rec_NonDefaultCtor> {name = 
"notDefCtor"}
+  // CHECK-NEXT: acc.serial 
firstprivate(@firstprivatization__ZTS14NonDefaultCtor -> %[[PRIVATE]] : 
!cir.ptr<!rec_NonDefaultCtor>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(dtor)
+  ;
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTOR]] : 
!cir.ptr<!rec_HasDtor>) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}
+  // CHECK-NEXT: acc.serial firstprivate(@firstprivatization__ZTS7HasDtor -> 
%[[PRIVATE]] : !cir.ptr<!rec_HasDtor>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel firstprivate(someInt, someFloat, noCopy, hasCopy, 
notDefCtor, dtor)
+  ;
+  // CHECK: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[SOMEINT]] : 
!cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[SOMEFLOAT]] : 
!cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPY]] : 
!cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = 
"noCopy"}
+  // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPY]] : 
!cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTOR]] : 
!cir.ptr<!rec_NonDefaultCtor>) -> !cir.ptr<!rec_NonDefaultCtor> {name = 
"notDefCtor"}
+  // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTOR]] : 
!cir.ptr<!rec_HasDtor>) -> !cir.ptr<!rec_HasDtor> {name = "dtor"}
+  // CHECK-NEXT: acc.parallel firstprivate(@firstprivatization__ZTSi -> 
%[[PRIVATE1]] : !cir.ptr<!s32i>,
+  // CHECK-SAME: @firstprivatization__ZTSf -> %[[PRIVATE2]] : 
!cir.ptr<!cir.float>,
+  // CHECK-SAME: @firstprivatization__ZTS15NoCopyConstruct -> %[[PRIVATE3]] : 
!cir.ptr<!rec_NoCopyConstruct>,
+  // CHECK-SAME: @firstprivatization__ZTS13CopyConstruct -> %[[PRIVATE4]] : 
!cir.ptr<!rec_CopyConstruct>,
+  // CHECK-SAME: @firstprivatization__ZTS14NonDefaultCtor -> %[[PRIVATE5]] : 
!cir.ptr<!rec_NonDefaultCtor>,
+  // CHECK-SAME: @firstprivatization__ZTS7HasDtor -> %[[PRIVATE6]] : 
!cir.ptr<!rec_HasDtor>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc serial firstprivate(someIntArr[1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}
+  // CHECK-NEXT: acc.serial firstprivate(@firstprivatization__ZTSA5_i -> 
%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(someFloatArr[1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}
+  // CHECK-NEXT: acc.parallel firstprivate(@firstprivatization__ZTSA5_f -> 
%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(noCopyArr[1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}
+  // CHECK-NEXT: acc.serial 
firstprivate(@firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(hasCopyArr[1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}
+  // CHECK-NEXT: acc.parallel 
firstprivate(@firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(notDefCtorArr[1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}
+  // CHECK-NEXT: acc.parallel 
firstprivate(@firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(dtorArr[1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}
+  // CHECK-NEXT: acc.parallel firstprivate(@firstprivatization__ZTSA5_7HasDtor 
-> %[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(someIntArr[1], someFloatArr[1], noCopyArr[1], 
hasCopyArr[1], notDefCtorArr[1], dtorArr[1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CONST]] : i64) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1]"}
+  // CHECK-NEXT: acc.serial firstprivate(@firstprivatization__ZTSA5_i -> 
%[[PRIVATE1]] : !cir.ptr<!cir.array<!s32i x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_f -> %[[PRIVATE2]] : 
!cir.ptr<!cir.array<!cir.float x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE3]] 
: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE4]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE5]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_7HasDtor -> %[[PRIVATE6]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel firstprivate(someIntArr[1:1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}
+  // CHECK-NEXT: acc.parallel firstprivate(@firstprivatization__ZTSA5_i -> 
%[[PRIVATE]] : !cir.ptr<!cir.array<!s32i x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(someFloatArr[1:1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}
+  // CHECK-NEXT: acc.serial firstprivate(@firstprivatization__ZTSA5_f -> 
%[[PRIVATE]] : !cir.ptr<!cir.array<!cir.float x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(noCopyArr[1:1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}
+  // CHECK-NEXT: acc.parallel 
firstprivate(@firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial firstprivate(hasCopyArr[1:1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}
+  // CHECK-NEXT: acc.serial 
firstprivate(@firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(notDefCtorArr[1:1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}
+  // CHECK-NEXT: acc.parallel 
firstprivate(@firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE]] : 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(dtorArr[1:1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}
+  // CHECK-NEXT: acc.parallel firstprivate(@firstprivatization__ZTSA5_7HasDtor 
-> %[[PRIVATE]] : !cir.ptr<!cir.array<!rec_HasDtor x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc parallel firstprivate(someIntArr[1:1], someFloatArr[1:1], 
noCopyArr[1:1], hasCopyArr[1:1], notDefCtorArr[1:1], dtorArr[1:1])
+  ;
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE1:.*]] = acc.firstprivate varPtr(%[[INTARR]] : 
!cir.ptr<!cir.array<!s32i x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!s32i x 5>> {name = "someIntArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE2:.*]] = acc.firstprivate varPtr(%[[FLOATARR]] : 
!cir.ptr<!cir.array<!cir.float x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!cir.float x 5>> {name = "someFloatArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE3:.*]] = acc.firstprivate varPtr(%[[NOCOPYARR]] : 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>> {name = "noCopyArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE4:.*]] = acc.firstprivate varPtr(%[[HASCOPYARR]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>> {name = "hasCopyArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE5:.*]] = acc.firstprivate varPtr(%[[NOTDEFCTORARR]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>> {name = "notDefCtorArr[1:1]"}
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i
+  // CHECK-NEXT: %[[ONE_CAST2:.*]] = builtin.unrealized_conversion_cast 
%[[ONE]] : !s32i to si32
+  // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0
+  // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1
+  // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) 
extent(%[[ONE_CAST2]] : si32) stride(%[[ONE_CONST2]] : i64) 
startIdx(%[[ZERO_CONST]] : i64)
+  // CHECK-NEXT: %[[PRIVATE6:.*]] = acc.firstprivate varPtr(%[[DTORARR]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>) bounds(%[[BOUNDS]]) -> 
!cir.ptr<!cir.array<!rec_HasDtor x 5>> {name = "dtorArr[1:1]"}
+  // CHECK-NEXT: acc.parallel firstprivate(@firstprivatization__ZTSA5_i -> 
%[[PRIVATE1]] : !cir.ptr<!cir.array<!s32i x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_f -> %[[PRIVATE2]] : 
!cir.ptr<!cir.array<!cir.float x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_15NoCopyConstruct -> %[[PRIVATE3]] 
: !cir.ptr<!cir.array<!rec_NoCopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_13CopyConstruct -> %[[PRIVATE4]] : 
!cir.ptr<!cir.array<!rec_CopyConstruct x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_14NonDefaultCtor -> %[[PRIVATE5]] 
: !cir.ptr<!cir.array<!rec_NonDefaultCtor x 5>>,
+  // CHECK-SAME: @firstprivatization__ZTSA5_7HasDtor -> %[[PRIVATE6]] : 
!cir.ptr<!cir.array<!rec_HasDtor x 5>>)
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+}

diff  --git a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp 
b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
index 0bf932ea62ceb..da45aca13e7f9 100644
--- a/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
+++ b/clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp
@@ -10,9 +10,6 @@ void HelloWorld(int *A, int *B, int *C, int N) {
 // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Declare 
Construct}}
 #pragma acc declare create(A)
 
-  // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: 
firstprivate}}
-#pragma acc parallel loop firstprivate(A)
-  for(int i = 0; i <5; ++i);
   // expected-error@+1{{ClangIR code gen Not Yet Implemented: OpenACC Clause: 
reduction}}
 #pragma acc parallel loop reduction(+:A)
   for(int i = 0; i <5; ++i);


        
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

[clang] 340fa3e - [OpenACC] Implement firstprivate lowering except init. (#153847)

Reply via email to