[OpenACC][CIR] Implement 'private' clause lowering. #151360

erichkeane · 2025-07-30T16:15:18Z

The private clause is the first with 'recipes', so a lot of infrastructure is included here, including some MLIR dialect changes that allow simple adding of a privatization. We'll likely get similar for firstprivate and reduction.

Also, we have quite a bit of infrastructure in clause lowering to make sure we have most cases we could think of covered.

At the moment, ONLY private is implemented, so all it requires is an 'init' segment (that doesn't call any copy operations), and potentially a 'destroy' segment. However, actually calling 'init' functions on each of the elements in them are not properly implemented, and will be in a followup patch.

This patch implements all of that, and adds tests in a way that will be useful for firstprivate as well.

The private clause is the first with 'recipes', so a lot of infrastructure is included here, including some MLIR dialect changes that allow simple adding of a privatization. We'll likely get similar for firstprivate and reduction. Also, we have quite a bit of infrastructure in clause lowering to make sure we have most cases we could think of covered. At the moment, ONLY private is implemented, so all it requires is an 'init' segment (that doesn't call any copy operations), and potentially a 'destroy' segment. This patch implements all of that, and adds tests in a way that will be useful for firstprivate as well.

llvmbot · 2025-07-30T16:15:53Z

@llvm/pr-subscribers-mlir

@llvm/pr-subscribers-clangir

Author: Erich Keane (erichkeane)

Changes

The private clause is the first with 'recipes', so a lot of infrastructure is included here, including some MLIR dialect changes that allow simple adding of a privatization. We'll likely get similar for firstprivate and reduction.

Also, we have quite a bit of infrastructure in clause lowering to make sure we have most cases we could think of covered.

At the moment, ONLY private is implemented, so all it requires is an 'init' segment (that doesn't call any copy operations), and potentially a 'destroy' segment.

This patch implements all of that, and adds tests in a way that will be useful for firstprivate as well.

Patch is 120.27 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/151360.diff

10 Files Affected:

(modified) clang/lib/CIR/CodeGen/CIRGenDecl.cpp (+4-2)
(modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+3-1)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp (+2-2)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+136)
(added) clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp (+522)
(added) clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp (+459)
(added) clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp (+459)
(modified) clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp (+2-2)
(modified) mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td (+11)
(modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+45)

diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index 6527fb5697f7c..5052e35f120cb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -24,7 +24,8 @@ using namespace clang;
 using namespace clang::CIRGen;
 
 CIRGenFunction::AutoVarEmission
-CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
+CIRGenFunction::emitAutoVarAlloca(const VarDecl &d,
+                                  mlir::OpBuilder::InsertPoint ip) {
   QualType ty = d.getType();
   if (ty.getAddressSpace() != LangAS::Default)
     cgm.errorNYI(d.getSourceRange(), "emitAutoVarAlloca: address space");
@@ -50,7 +51,8 @@ CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
   // A normal fixed sized variable becomes an alloca in the entry block,
   mlir::Type allocaTy = convertTypeForMem(ty);
   // Create the temp alloca and declare variable using it.
-  address = createTempAlloca(allocaTy, alignment, loc, d.getName());
+  address = createTempAlloca(allocaTy, alignment, loc, d.getName(),
+                             /*arraySize=*/nullptr, /*alloca=*/nullptr, ip);
   declare(address.getPointer(), &d, ty, getLoc(d.getSourceRange()), alignment);
 
   emission.Addr = address;
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 603f75078c519..830f9e8d72b68 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -858,7 +858,8 @@ class CIRGenFunction : public CIRGenTypeCache {
 
   Address emitArrayToPointerDecay(const Expr *array);
 
-  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d);
+  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d,
+                                    mlir::OpBuilder::InsertPoint ip = {});
 
   /// Emit code and set up symbol table for a variable declaration with auto,
   /// register, or no storage class specifier. These turn into simple stack
@@ -1375,6 +1376,7 @@ class CIRGenFunction : public CIRGenTypeCache {
     mlir::Location beginLoc;
     mlir::Value varValue;
     std::string name;
+    QualType baseType;
     llvm::SmallVector<mlir::Value> bounds;
   };
   // Gets the collection of info required to lower and OpenACC clause or cache
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
index 49ff1249827a4..32095cb687e88 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
@@ -119,7 +119,7 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
 
   if (const auto *memExpr = dyn_cast<MemberExpr>(curVarExpr))
     return {exprLoc, emitMemberExpr(memExpr).getPointer(), exprString,
-            std::move(bounds)};
+            curVarExpr->getType(), std::move(bounds)};
 
   // Sema has made sure that only 4 types of things can get here, array
   // subscript, array section, member expr, or DRE to a var decl (or the
@@ -127,5 +127,5 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
   // right.
   const auto *dre = cast<DeclRefExpr>(curVarExpr);
   return {exprLoc, emitDeclRefLValue(dre).getPointer(), exprString,
-          std::move(bounds)};
+          curVarExpr->getType(), std::move(bounds)};
 }
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index e45d3b8f4aa82..4e6fe41afcdd9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -12,6 +12,7 @@
 
 #include <type_traits>
 
+#include "CIRGenCXXABI.h"
 #include "CIRGenFunction.h"
 
 #include "clang/AST/ExprCXX.h"
@@ -355,6 +356,110 @@ class OpenACCClauseCIREmitter final
     }
   }
 
+  template <typename RecipeTy>
+  RecipeTy getOrCreateRecipe(ASTContext &astCtx, const Expr *varRef,
+                             DeclContext *dc, QualType baseType,
+                             mlir::Value mainOp) {
+    mlir::ModuleOp mod =
+        builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();
+
+    std::string recipeName;
+    {
+      llvm::raw_string_ostream stream(recipeName);
+      if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        stream << "privatization_";
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::FirstprivateRecipeOp>) {
+        stream << "firstprivatization_";
+
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::ReductionRecipeOp>) {
+        stream << "reduction_";
+        // We don't have the reduction operation here well enough to know how to
+        // spell this correctly (+ == 'add', etc), so when we implement
+        // 'reduction' we have to do that here.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpeNACC reduction recipe creation");
+      } else {
+        static_assert(!sizeof(RecipeTy), "Unknown Recipe op kind");
+      }
+
+      MangleContext &mc = cgf.cgm.getCXXABI().getMangleContext();
+      mc.mangleCanonicalTypeName(baseType, stream);
+    }
+
+    if (auto recipe = mod.lookupSymbol<RecipeTy>(recipeName))
+      return recipe;
+
+    mlir::Location loc = cgf.cgm.getLoc(varRef->getBeginLoc());
+    mlir::Location locEnd = cgf.cgm.getLoc(varRef->getEndLoc());
+
+    mlir::OpBuilder modBuilder(mod.getBodyRegion());
+    auto recipe =
+        modBuilder.create<RecipeTy>(loc, recipeName, mainOp.getType());
+
+    // Magic-up a var-decl so we can use normal init/destruction operations for
+    // a variable declaration.
+    VarDecl &tempDecl = *VarDecl::Create(
+        astCtx, dc, varRef->getBeginLoc(), varRef->getBeginLoc(),
+        &astCtx.Idents.get("openacc.private.init"), baseType,
+        astCtx.getTrivialTypeSourceInfo(baseType), SC_Auto);
+    CIRGenFunction::AutoVarEmission tempDeclEmission{
+        CIRGenFunction::AutoVarEmission::invalid()};
+
+    // Init section.
+    {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      builder.createBlock(&recipe.getInitRegion(), recipe.getInitRegion().end(),
+                          argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
+
+      if constexpr (!std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        // We have only implemented 'init' for private, so make this NYI until
+        // we have explicitly implemented everything.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpenACC non-private recipe init");
+      }
+
+      tempDeclEmission =
+          cgf.emitAutoVarAlloca(tempDecl, builder.saveInsertionPoint());
+      cgf.emitAutoVarInit(tempDeclEmission);
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    // Copy section.
+    if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp> ||
+                  std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
+      // TODO: OpenACC: 'private' doesn't emit this, but for the other two we
+      // have to figure out what 'copy' means here.
+      cgf.cgm.errorNYI(varRef->getSourceRange(),
+                       "OpenACC record type privatization copy section");
+    }
+
+    // Destroy section (doesn't currently exist).
+    if (tempDecl.needsDestruction(cgf.getContext())) {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      mlir::Block *block = builder.createBlock(&recipe.getDestroyRegion(),
+                                               recipe.getDestroyRegion().end(),
+                                               argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getDestroyRegion().back());
+
+      mlir::Type elementTy =
+          mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
+      Address addr{block->getArgument(0), elementTy,
+                   cgf.getContext().getDeclAlign(&tempDecl)};
+      cgf.emitDestroy(addr, baseType,
+                      cgf.getDestroyer(QualType::DK_cxx_destructor));
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    return recipe;
+  }
+
 public:
   OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
                           CIRGen::CIRGenBuilderTy &builder,
@@ -971,6 +1076,37 @@ class OpenACCClauseCIREmitter final
       llvm_unreachable("Unknown construct kind in VisitAttachClause");
     }
   }
+
+  void VisitPrivateClause(const OpenACCPrivateClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
+                               mlir::acc::LoopOp>) {
+      for (const Expr *var : clause.getVarList()) {
+        CIRGenFunction::OpenACCDataOperandInfo opInfo =
+            cgf.getOpenACCDataOperandInfo(var);
+        auto privateOp = builder.create<mlir::acc::PrivateOp>(
+            opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
+            /*implicit=*/false, opInfo.name, opInfo.bounds);
+        privateOp.setDataClause(mlir::acc::DataClause::acc_private);
+
+        {
+          mlir::OpBuilder::InsertionGuard guardCase(builder);
+          auto recipe = getOrCreateRecipe<mlir::acc::PrivateRecipeOp>(
+              cgf.getContext(), var, Decl::castToDeclContext(cgf.curFuncDecl),
+              opInfo.baseType, privateOp.getResult());
+          // TODO: OpenACC: The dialect is going to change in the near future to
+          // have these be on a different operation, so when that changes, we
+          // probably need to change these here.
+          operation.addPrivatization(builder.getContext(), privateOp, recipe);
+        }
+      }
+    } else if constexpr (isCombinedType<OpTy>) {
+      // Despite this being valid on ParallelOp or SerialOp, combined type
+      // applies to the 'loop'.
+      applyToLoopOp(clause);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitPrivateClause");
+    }
+  }
 };
 
 template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
new file mode 100644
index 0000000000000..3306c55b4a8e7
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
@@ -0,0 +1,522 @@
+// 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.private.recipe @privatization__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: } 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.private.recipe @privatization__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: 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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: } 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.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+
+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 private(someInt)
+  for(int i = 0; i < 5; ++i);
+  // CHECK: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTSi -> %[[PRIVATE]] : !cir.ptr<!s32i>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(someFloat)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTSf -> %[[PRIVATE]] : !cir.ptr<!cir.float>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel loop private(noCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTS15NoCopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(hasCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTS13CopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEX...
[truncated]

llvmbot · 2025-07-30T16:15:53Z

@llvm/pr-subscribers-mlir-openacc

Author: Erich Keane (erichkeane)

Changes

The private clause is the first with 'recipes', so a lot of infrastructure is included here, including some MLIR dialect changes that allow simple adding of a privatization. We'll likely get similar for firstprivate and reduction.

Also, we have quite a bit of infrastructure in clause lowering to make sure we have most cases we could think of covered.

At the moment, ONLY private is implemented, so all it requires is an 'init' segment (that doesn't call any copy operations), and potentially a 'destroy' segment.

This patch implements all of that, and adds tests in a way that will be useful for firstprivate as well.

Patch is 120.27 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/151360.diff

10 Files Affected:

(modified) clang/lib/CIR/CodeGen/CIRGenDecl.cpp (+4-2)
(modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+3-1)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp (+2-2)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+136)
(added) clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp (+522)
(added) clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp (+459)
(added) clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp (+459)
(modified) clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp (+2-2)
(modified) mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td (+11)
(modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+45)

diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index 6527fb5697f7c..5052e35f120cb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -24,7 +24,8 @@ using namespace clang;
 using namespace clang::CIRGen;
 
 CIRGenFunction::AutoVarEmission
-CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
+CIRGenFunction::emitAutoVarAlloca(const VarDecl &d,
+                                  mlir::OpBuilder::InsertPoint ip) {
   QualType ty = d.getType();
   if (ty.getAddressSpace() != LangAS::Default)
     cgm.errorNYI(d.getSourceRange(), "emitAutoVarAlloca: address space");
@@ -50,7 +51,8 @@ CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
   // A normal fixed sized variable becomes an alloca in the entry block,
   mlir::Type allocaTy = convertTypeForMem(ty);
   // Create the temp alloca and declare variable using it.
-  address = createTempAlloca(allocaTy, alignment, loc, d.getName());
+  address = createTempAlloca(allocaTy, alignment, loc, d.getName(),
+                             /*arraySize=*/nullptr, /*alloca=*/nullptr, ip);
   declare(address.getPointer(), &d, ty, getLoc(d.getSourceRange()), alignment);
 
   emission.Addr = address;
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 603f75078c519..830f9e8d72b68 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -858,7 +858,8 @@ class CIRGenFunction : public CIRGenTypeCache {
 
   Address emitArrayToPointerDecay(const Expr *array);
 
-  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d);
+  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d,
+                                    mlir::OpBuilder::InsertPoint ip = {});
 
   /// Emit code and set up symbol table for a variable declaration with auto,
   /// register, or no storage class specifier. These turn into simple stack
@@ -1375,6 +1376,7 @@ class CIRGenFunction : public CIRGenTypeCache {
     mlir::Location beginLoc;
     mlir::Value varValue;
     std::string name;
+    QualType baseType;
     llvm::SmallVector<mlir::Value> bounds;
   };
   // Gets the collection of info required to lower and OpenACC clause or cache
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
index 49ff1249827a4..32095cb687e88 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
@@ -119,7 +119,7 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
 
   if (const auto *memExpr = dyn_cast<MemberExpr>(curVarExpr))
     return {exprLoc, emitMemberExpr(memExpr).getPointer(), exprString,
-            std::move(bounds)};
+            curVarExpr->getType(), std::move(bounds)};
 
   // Sema has made sure that only 4 types of things can get here, array
   // subscript, array section, member expr, or DRE to a var decl (or the
@@ -127,5 +127,5 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
   // right.
   const auto *dre = cast<DeclRefExpr>(curVarExpr);
   return {exprLoc, emitDeclRefLValue(dre).getPointer(), exprString,
-          std::move(bounds)};
+          curVarExpr->getType(), std::move(bounds)};
 }
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index e45d3b8f4aa82..4e6fe41afcdd9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -12,6 +12,7 @@
 
 #include <type_traits>
 
+#include "CIRGenCXXABI.h"
 #include "CIRGenFunction.h"
 
 #include "clang/AST/ExprCXX.h"
@@ -355,6 +356,110 @@ class OpenACCClauseCIREmitter final
     }
   }
 
+  template <typename RecipeTy>
+  RecipeTy getOrCreateRecipe(ASTContext &astCtx, const Expr *varRef,
+                             DeclContext *dc, QualType baseType,
+                             mlir::Value mainOp) {
+    mlir::ModuleOp mod =
+        builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();
+
+    std::string recipeName;
+    {
+      llvm::raw_string_ostream stream(recipeName);
+      if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        stream << "privatization_";
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::FirstprivateRecipeOp>) {
+        stream << "firstprivatization_";
+
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::ReductionRecipeOp>) {
+        stream << "reduction_";
+        // We don't have the reduction operation here well enough to know how to
+        // spell this correctly (+ == 'add', etc), so when we implement
+        // 'reduction' we have to do that here.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpeNACC reduction recipe creation");
+      } else {
+        static_assert(!sizeof(RecipeTy), "Unknown Recipe op kind");
+      }
+
+      MangleContext &mc = cgf.cgm.getCXXABI().getMangleContext();
+      mc.mangleCanonicalTypeName(baseType, stream);
+    }
+
+    if (auto recipe = mod.lookupSymbol<RecipeTy>(recipeName))
+      return recipe;
+
+    mlir::Location loc = cgf.cgm.getLoc(varRef->getBeginLoc());
+    mlir::Location locEnd = cgf.cgm.getLoc(varRef->getEndLoc());
+
+    mlir::OpBuilder modBuilder(mod.getBodyRegion());
+    auto recipe =
+        modBuilder.create<RecipeTy>(loc, recipeName, mainOp.getType());
+
+    // Magic-up a var-decl so we can use normal init/destruction operations for
+    // a variable declaration.
+    VarDecl &tempDecl = *VarDecl::Create(
+        astCtx, dc, varRef->getBeginLoc(), varRef->getBeginLoc(),
+        &astCtx.Idents.get("openacc.private.init"), baseType,
+        astCtx.getTrivialTypeSourceInfo(baseType), SC_Auto);
+    CIRGenFunction::AutoVarEmission tempDeclEmission{
+        CIRGenFunction::AutoVarEmission::invalid()};
+
+    // Init section.
+    {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      builder.createBlock(&recipe.getInitRegion(), recipe.getInitRegion().end(),
+                          argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
+
+      if constexpr (!std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        // We have only implemented 'init' for private, so make this NYI until
+        // we have explicitly implemented everything.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpenACC non-private recipe init");
+      }
+
+      tempDeclEmission =
+          cgf.emitAutoVarAlloca(tempDecl, builder.saveInsertionPoint());
+      cgf.emitAutoVarInit(tempDeclEmission);
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    // Copy section.
+    if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp> ||
+                  std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
+      // TODO: OpenACC: 'private' doesn't emit this, but for the other two we
+      // have to figure out what 'copy' means here.
+      cgf.cgm.errorNYI(varRef->getSourceRange(),
+                       "OpenACC record type privatization copy section");
+    }
+
+    // Destroy section (doesn't currently exist).
+    if (tempDecl.needsDestruction(cgf.getContext())) {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      mlir::Block *block = builder.createBlock(&recipe.getDestroyRegion(),
+                                               recipe.getDestroyRegion().end(),
+                                               argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getDestroyRegion().back());
+
+      mlir::Type elementTy =
+          mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
+      Address addr{block->getArgument(0), elementTy,
+                   cgf.getContext().getDeclAlign(&tempDecl)};
+      cgf.emitDestroy(addr, baseType,
+                      cgf.getDestroyer(QualType::DK_cxx_destructor));
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    return recipe;
+  }
+
 public:
   OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
                           CIRGen::CIRGenBuilderTy &builder,
@@ -971,6 +1076,37 @@ class OpenACCClauseCIREmitter final
       llvm_unreachable("Unknown construct kind in VisitAttachClause");
     }
   }
+
+  void VisitPrivateClause(const OpenACCPrivateClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
+                               mlir::acc::LoopOp>) {
+      for (const Expr *var : clause.getVarList()) {
+        CIRGenFunction::OpenACCDataOperandInfo opInfo =
+            cgf.getOpenACCDataOperandInfo(var);
+        auto privateOp = builder.create<mlir::acc::PrivateOp>(
+            opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
+            /*implicit=*/false, opInfo.name, opInfo.bounds);
+        privateOp.setDataClause(mlir::acc::DataClause::acc_private);
+
+        {
+          mlir::OpBuilder::InsertionGuard guardCase(builder);
+          auto recipe = getOrCreateRecipe<mlir::acc::PrivateRecipeOp>(
+              cgf.getContext(), var, Decl::castToDeclContext(cgf.curFuncDecl),
+              opInfo.baseType, privateOp.getResult());
+          // TODO: OpenACC: The dialect is going to change in the near future to
+          // have these be on a different operation, so when that changes, we
+          // probably need to change these here.
+          operation.addPrivatization(builder.getContext(), privateOp, recipe);
+        }
+      }
+    } else if constexpr (isCombinedType<OpTy>) {
+      // Despite this being valid on ParallelOp or SerialOp, combined type
+      // applies to the 'loop'.
+      applyToLoopOp(clause);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitPrivateClause");
+    }
+  }
 };
 
 template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
new file mode 100644
index 0000000000000..3306c55b4a8e7
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
@@ -0,0 +1,522 @@
+// 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.private.recipe @privatization__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: } 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.private.recipe @privatization__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: 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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: } 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.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+
+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 private(someInt)
+  for(int i = 0; i < 5; ++i);
+  // CHECK: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTSi -> %[[PRIVATE]] : !cir.ptr<!s32i>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(someFloat)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTSf -> %[[PRIVATE]] : !cir.ptr<!cir.float>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel loop private(noCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTS15NoCopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(hasCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTS13CopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEX...
[truncated]

llvmbot · 2025-07-30T16:15:54Z

@llvm/pr-subscribers-openacc

Author: Erich Keane (erichkeane)

Changes

The private clause is the first with 'recipes', so a lot of infrastructure is included here, including some MLIR dialect changes that allow simple adding of a privatization. We'll likely get similar for firstprivate and reduction.

Also, we have quite a bit of infrastructure in clause lowering to make sure we have most cases we could think of covered.

At the moment, ONLY private is implemented, so all it requires is an 'init' segment (that doesn't call any copy operations), and potentially a 'destroy' segment.

This patch implements all of that, and adds tests in a way that will be useful for firstprivate as well.

Patch is 120.27 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/151360.diff

10 Files Affected:

(modified) clang/lib/CIR/CodeGen/CIRGenDecl.cpp (+4-2)
(modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+3-1)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp (+2-2)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+136)
(added) clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp (+522)
(added) clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp (+459)
(added) clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp (+459)
(modified) clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp (+2-2)
(modified) mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td (+11)
(modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+45)

diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index 6527fb5697f7c..5052e35f120cb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -24,7 +24,8 @@ using namespace clang;
 using namespace clang::CIRGen;
 
 CIRGenFunction::AutoVarEmission
-CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
+CIRGenFunction::emitAutoVarAlloca(const VarDecl &d,
+                                  mlir::OpBuilder::InsertPoint ip) {
   QualType ty = d.getType();
   if (ty.getAddressSpace() != LangAS::Default)
     cgm.errorNYI(d.getSourceRange(), "emitAutoVarAlloca: address space");
@@ -50,7 +51,8 @@ CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
   // A normal fixed sized variable becomes an alloca in the entry block,
   mlir::Type allocaTy = convertTypeForMem(ty);
   // Create the temp alloca and declare variable using it.
-  address = createTempAlloca(allocaTy, alignment, loc, d.getName());
+  address = createTempAlloca(allocaTy, alignment, loc, d.getName(),
+                             /*arraySize=*/nullptr, /*alloca=*/nullptr, ip);
   declare(address.getPointer(), &d, ty, getLoc(d.getSourceRange()), alignment);
 
   emission.Addr = address;
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 603f75078c519..830f9e8d72b68 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -858,7 +858,8 @@ class CIRGenFunction : public CIRGenTypeCache {
 
   Address emitArrayToPointerDecay(const Expr *array);
 
-  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d);
+  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d,
+                                    mlir::OpBuilder::InsertPoint ip = {});
 
   /// Emit code and set up symbol table for a variable declaration with auto,
   /// register, or no storage class specifier. These turn into simple stack
@@ -1375,6 +1376,7 @@ class CIRGenFunction : public CIRGenTypeCache {
     mlir::Location beginLoc;
     mlir::Value varValue;
     std::string name;
+    QualType baseType;
     llvm::SmallVector<mlir::Value> bounds;
   };
   // Gets the collection of info required to lower and OpenACC clause or cache
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
index 49ff1249827a4..32095cb687e88 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
@@ -119,7 +119,7 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
 
   if (const auto *memExpr = dyn_cast<MemberExpr>(curVarExpr))
     return {exprLoc, emitMemberExpr(memExpr).getPointer(), exprString,
-            std::move(bounds)};
+            curVarExpr->getType(), std::move(bounds)};
 
   // Sema has made sure that only 4 types of things can get here, array
   // subscript, array section, member expr, or DRE to a var decl (or the
@@ -127,5 +127,5 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
   // right.
   const auto *dre = cast<DeclRefExpr>(curVarExpr);
   return {exprLoc, emitDeclRefLValue(dre).getPointer(), exprString,
-          std::move(bounds)};
+          curVarExpr->getType(), std::move(bounds)};
 }
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index e45d3b8f4aa82..4e6fe41afcdd9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -12,6 +12,7 @@
 
 #include <type_traits>
 
+#include "CIRGenCXXABI.h"
 #include "CIRGenFunction.h"
 
 #include "clang/AST/ExprCXX.h"
@@ -355,6 +356,110 @@ class OpenACCClauseCIREmitter final
     }
   }
 
+  template <typename RecipeTy>
+  RecipeTy getOrCreateRecipe(ASTContext &astCtx, const Expr *varRef,
+                             DeclContext *dc, QualType baseType,
+                             mlir::Value mainOp) {
+    mlir::ModuleOp mod =
+        builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();
+
+    std::string recipeName;
+    {
+      llvm::raw_string_ostream stream(recipeName);
+      if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        stream << "privatization_";
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::FirstprivateRecipeOp>) {
+        stream << "firstprivatization_";
+
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::ReductionRecipeOp>) {
+        stream << "reduction_";
+        // We don't have the reduction operation here well enough to know how to
+        // spell this correctly (+ == 'add', etc), so when we implement
+        // 'reduction' we have to do that here.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpeNACC reduction recipe creation");
+      } else {
+        static_assert(!sizeof(RecipeTy), "Unknown Recipe op kind");
+      }
+
+      MangleContext &mc = cgf.cgm.getCXXABI().getMangleContext();
+      mc.mangleCanonicalTypeName(baseType, stream);
+    }
+
+    if (auto recipe = mod.lookupSymbol<RecipeTy>(recipeName))
+      return recipe;
+
+    mlir::Location loc = cgf.cgm.getLoc(varRef->getBeginLoc());
+    mlir::Location locEnd = cgf.cgm.getLoc(varRef->getEndLoc());
+
+    mlir::OpBuilder modBuilder(mod.getBodyRegion());
+    auto recipe =
+        modBuilder.create<RecipeTy>(loc, recipeName, mainOp.getType());
+
+    // Magic-up a var-decl so we can use normal init/destruction operations for
+    // a variable declaration.
+    VarDecl &tempDecl = *VarDecl::Create(
+        astCtx, dc, varRef->getBeginLoc(), varRef->getBeginLoc(),
+        &astCtx.Idents.get("openacc.private.init"), baseType,
+        astCtx.getTrivialTypeSourceInfo(baseType), SC_Auto);
+    CIRGenFunction::AutoVarEmission tempDeclEmission{
+        CIRGenFunction::AutoVarEmission::invalid()};
+
+    // Init section.
+    {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      builder.createBlock(&recipe.getInitRegion(), recipe.getInitRegion().end(),
+                          argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
+
+      if constexpr (!std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        // We have only implemented 'init' for private, so make this NYI until
+        // we have explicitly implemented everything.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpenACC non-private recipe init");
+      }
+
+      tempDeclEmission =
+          cgf.emitAutoVarAlloca(tempDecl, builder.saveInsertionPoint());
+      cgf.emitAutoVarInit(tempDeclEmission);
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    // Copy section.
+    if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp> ||
+                  std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
+      // TODO: OpenACC: 'private' doesn't emit this, but for the other two we
+      // have to figure out what 'copy' means here.
+      cgf.cgm.errorNYI(varRef->getSourceRange(),
+                       "OpenACC record type privatization copy section");
+    }
+
+    // Destroy section (doesn't currently exist).
+    if (tempDecl.needsDestruction(cgf.getContext())) {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      mlir::Block *block = builder.createBlock(&recipe.getDestroyRegion(),
+                                               recipe.getDestroyRegion().end(),
+                                               argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getDestroyRegion().back());
+
+      mlir::Type elementTy =
+          mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
+      Address addr{block->getArgument(0), elementTy,
+                   cgf.getContext().getDeclAlign(&tempDecl)};
+      cgf.emitDestroy(addr, baseType,
+                      cgf.getDestroyer(QualType::DK_cxx_destructor));
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    return recipe;
+  }
+
 public:
   OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
                           CIRGen::CIRGenBuilderTy &builder,
@@ -971,6 +1076,37 @@ class OpenACCClauseCIREmitter final
       llvm_unreachable("Unknown construct kind in VisitAttachClause");
     }
   }
+
+  void VisitPrivateClause(const OpenACCPrivateClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
+                               mlir::acc::LoopOp>) {
+      for (const Expr *var : clause.getVarList()) {
+        CIRGenFunction::OpenACCDataOperandInfo opInfo =
+            cgf.getOpenACCDataOperandInfo(var);
+        auto privateOp = builder.create<mlir::acc::PrivateOp>(
+            opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
+            /*implicit=*/false, opInfo.name, opInfo.bounds);
+        privateOp.setDataClause(mlir::acc::DataClause::acc_private);
+
+        {
+          mlir::OpBuilder::InsertionGuard guardCase(builder);
+          auto recipe = getOrCreateRecipe<mlir::acc::PrivateRecipeOp>(
+              cgf.getContext(), var, Decl::castToDeclContext(cgf.curFuncDecl),
+              opInfo.baseType, privateOp.getResult());
+          // TODO: OpenACC: The dialect is going to change in the near future to
+          // have these be on a different operation, so when that changes, we
+          // probably need to change these here.
+          operation.addPrivatization(builder.getContext(), privateOp, recipe);
+        }
+      }
+    } else if constexpr (isCombinedType<OpTy>) {
+      // Despite this being valid on ParallelOp or SerialOp, combined type
+      // applies to the 'loop'.
+      applyToLoopOp(clause);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitPrivateClause");
+    }
+  }
 };
 
 template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
new file mode 100644
index 0000000000000..3306c55b4a8e7
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
@@ -0,0 +1,522 @@
+// 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.private.recipe @privatization__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: } 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.private.recipe @privatization__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: 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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: } 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.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+
+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 private(someInt)
+  for(int i = 0; i < 5; ++i);
+  // CHECK: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTSi -> %[[PRIVATE]] : !cir.ptr<!s32i>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(someFloat)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTSf -> %[[PRIVATE]] : !cir.ptr<!cir.float>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel loop private(noCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTS15NoCopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(hasCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTS13CopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEX...
[truncated]

llvmbot · 2025-07-30T16:15:54Z

@llvm/pr-subscribers-clang

Author: Erich Keane (erichkeane)

Changes

The private clause is the first with 'recipes', so a lot of infrastructure is included here, including some MLIR dialect changes that allow simple adding of a privatization. We'll likely get similar for firstprivate and reduction.

Also, we have quite a bit of infrastructure in clause lowering to make sure we have most cases we could think of covered.

At the moment, ONLY private is implemented, so all it requires is an 'init' segment (that doesn't call any copy operations), and potentially a 'destroy' segment.

This patch implements all of that, and adds tests in a way that will be useful for firstprivate as well.

Patch is 120.27 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/151360.diff

10 Files Affected:

(modified) clang/lib/CIR/CodeGen/CIRGenDecl.cpp (+4-2)
(modified) clang/lib/CIR/CodeGen/CIRGenFunction.h (+3-1)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp (+2-2)
(modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+136)
(added) clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp (+522)
(added) clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp (+459)
(added) clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp (+459)
(modified) clang/test/CIR/CodeGenOpenACC/openacc-not-implemented.cpp (+2-2)
(modified) mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td (+11)
(modified) mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp (+45)

diff --git a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
index 6527fb5697f7c..5052e35f120cb 100644
--- a/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenDecl.cpp
@@ -24,7 +24,8 @@ using namespace clang;
 using namespace clang::CIRGen;
 
 CIRGenFunction::AutoVarEmission
-CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
+CIRGenFunction::emitAutoVarAlloca(const VarDecl &d,
+                                  mlir::OpBuilder::InsertPoint ip) {
   QualType ty = d.getType();
   if (ty.getAddressSpace() != LangAS::Default)
     cgm.errorNYI(d.getSourceRange(), "emitAutoVarAlloca: address space");
@@ -50,7 +51,8 @@ CIRGenFunction::emitAutoVarAlloca(const VarDecl &d) {
   // A normal fixed sized variable becomes an alloca in the entry block,
   mlir::Type allocaTy = convertTypeForMem(ty);
   // Create the temp alloca and declare variable using it.
-  address = createTempAlloca(allocaTy, alignment, loc, d.getName());
+  address = createTempAlloca(allocaTy, alignment, loc, d.getName(),
+                             /*arraySize=*/nullptr, /*alloca=*/nullptr, ip);
   declare(address.getPointer(), &d, ty, getLoc(d.getSourceRange()), alignment);
 
   emission.Addr = address;
diff --git a/clang/lib/CIR/CodeGen/CIRGenFunction.h b/clang/lib/CIR/CodeGen/CIRGenFunction.h
index 603f75078c519..830f9e8d72b68 100644
--- a/clang/lib/CIR/CodeGen/CIRGenFunction.h
+++ b/clang/lib/CIR/CodeGen/CIRGenFunction.h
@@ -858,7 +858,8 @@ class CIRGenFunction : public CIRGenTypeCache {
 
   Address emitArrayToPointerDecay(const Expr *array);
 
-  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d);
+  AutoVarEmission emitAutoVarAlloca(const clang::VarDecl &d,
+                                    mlir::OpBuilder::InsertPoint ip = {});
 
   /// Emit code and set up symbol table for a variable declaration with auto,
   /// register, or no storage class specifier. These turn into simple stack
@@ -1375,6 +1376,7 @@ class CIRGenFunction : public CIRGenTypeCache {
     mlir::Location beginLoc;
     mlir::Value varValue;
     std::string name;
+    QualType baseType;
     llvm::SmallVector<mlir::Value> bounds;
   };
   // Gets the collection of info required to lower and OpenACC clause or cache
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
index 49ff1249827a4..32095cb687e88 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp
@@ -119,7 +119,7 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
 
   if (const auto *memExpr = dyn_cast<MemberExpr>(curVarExpr))
     return {exprLoc, emitMemberExpr(memExpr).getPointer(), exprString,
-            std::move(bounds)};
+            curVarExpr->getType(), std::move(bounds)};
 
   // Sema has made sure that only 4 types of things can get here, array
   // subscript, array section, member expr, or DRE to a var decl (or the
@@ -127,5 +127,5 @@ CIRGenFunction::getOpenACCDataOperandInfo(const Expr *e) {
   // right.
   const auto *dre = cast<DeclRefExpr>(curVarExpr);
   return {exprLoc, emitDeclRefLValue(dre).getPointer(), exprString,
-          std::move(bounds)};
+          curVarExpr->getType(), std::move(bounds)};
 }
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index e45d3b8f4aa82..4e6fe41afcdd9 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -12,6 +12,7 @@
 
 #include <type_traits>
 
+#include "CIRGenCXXABI.h"
 #include "CIRGenFunction.h"
 
 #include "clang/AST/ExprCXX.h"
@@ -355,6 +356,110 @@ class OpenACCClauseCIREmitter final
     }
   }
 
+  template <typename RecipeTy>
+  RecipeTy getOrCreateRecipe(ASTContext &astCtx, const Expr *varRef,
+                             DeclContext *dc, QualType baseType,
+                             mlir::Value mainOp) {
+    mlir::ModuleOp mod =
+        builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();
+
+    std::string recipeName;
+    {
+      llvm::raw_string_ostream stream(recipeName);
+      if constexpr (std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        stream << "privatization_";
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::FirstprivateRecipeOp>) {
+        stream << "firstprivatization_";
+
+      } else if constexpr (std::is_same_v<RecipeTy,
+                                          mlir::acc::ReductionRecipeOp>) {
+        stream << "reduction_";
+        // We don't have the reduction operation here well enough to know how to
+        // spell this correctly (+ == 'add', etc), so when we implement
+        // 'reduction' we have to do that here.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpeNACC reduction recipe creation");
+      } else {
+        static_assert(!sizeof(RecipeTy), "Unknown Recipe op kind");
+      }
+
+      MangleContext &mc = cgf.cgm.getCXXABI().getMangleContext();
+      mc.mangleCanonicalTypeName(baseType, stream);
+    }
+
+    if (auto recipe = mod.lookupSymbol<RecipeTy>(recipeName))
+      return recipe;
+
+    mlir::Location loc = cgf.cgm.getLoc(varRef->getBeginLoc());
+    mlir::Location locEnd = cgf.cgm.getLoc(varRef->getEndLoc());
+
+    mlir::OpBuilder modBuilder(mod.getBodyRegion());
+    auto recipe =
+        modBuilder.create<RecipeTy>(loc, recipeName, mainOp.getType());
+
+    // Magic-up a var-decl so we can use normal init/destruction operations for
+    // a variable declaration.
+    VarDecl &tempDecl = *VarDecl::Create(
+        astCtx, dc, varRef->getBeginLoc(), varRef->getBeginLoc(),
+        &astCtx.Idents.get("openacc.private.init"), baseType,
+        astCtx.getTrivialTypeSourceInfo(baseType), SC_Auto);
+    CIRGenFunction::AutoVarEmission tempDeclEmission{
+        CIRGenFunction::AutoVarEmission::invalid()};
+
+    // Init section.
+    {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      builder.createBlock(&recipe.getInitRegion(), recipe.getInitRegion().end(),
+                          argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getInitRegion().back());
+
+      if constexpr (!std::is_same_v<RecipeTy, mlir::acc::PrivateRecipeOp>) {
+        // We have only implemented 'init' for private, so make this NYI until
+        // we have explicitly implemented everything.
+        cgf.cgm.errorNYI(varRef->getSourceRange(),
+                         "OpenACC non-private recipe init");
+      }
+
+      tempDeclEmission =
+          cgf.emitAutoVarAlloca(tempDecl, builder.saveInsertionPoint());
+      cgf.emitAutoVarInit(tempDeclEmission);
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    // Copy section.
+    if constexpr (std::is_same_v<RecipeTy, mlir::acc::FirstprivateRecipeOp> ||
+                  std::is_same_v<RecipeTy, mlir::acc::ReductionRecipeOp>) {
+      // TODO: OpenACC: 'private' doesn't emit this, but for the other two we
+      // have to figure out what 'copy' means here.
+      cgf.cgm.errorNYI(varRef->getSourceRange(),
+                       "OpenACC record type privatization copy section");
+    }
+
+    // Destroy section (doesn't currently exist).
+    if (tempDecl.needsDestruction(cgf.getContext())) {
+      llvm::SmallVector<mlir::Type> argsTys{mainOp.getType()};
+      llvm::SmallVector<mlir::Location> argsLocs{loc};
+      mlir::Block *block = builder.createBlock(&recipe.getDestroyRegion(),
+                                               recipe.getDestroyRegion().end(),
+                                               argsTys, argsLocs);
+      builder.setInsertionPointToEnd(&recipe.getDestroyRegion().back());
+
+      mlir::Type elementTy =
+          mlir::cast<cir::PointerType>(mainOp.getType()).getPointee();
+      Address addr{block->getArgument(0), elementTy,
+                   cgf.getContext().getDeclAlign(&tempDecl)};
+      cgf.emitDestroy(addr, baseType,
+                      cgf.getDestroyer(QualType::DK_cxx_destructor));
+
+      builder.create<mlir::acc::YieldOp>(locEnd);
+    }
+
+    return recipe;
+  }
+
 public:
   OpenACCClauseCIREmitter(OpTy &operation, CIRGen::CIRGenFunction &cgf,
                           CIRGen::CIRGenBuilderTy &builder,
@@ -971,6 +1076,37 @@ class OpenACCClauseCIREmitter final
       llvm_unreachable("Unknown construct kind in VisitAttachClause");
     }
   }
+
+  void VisitPrivateClause(const OpenACCPrivateClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::ParallelOp, mlir::acc::SerialOp,
+                               mlir::acc::LoopOp>) {
+      for (const Expr *var : clause.getVarList()) {
+        CIRGenFunction::OpenACCDataOperandInfo opInfo =
+            cgf.getOpenACCDataOperandInfo(var);
+        auto privateOp = builder.create<mlir::acc::PrivateOp>(
+            opInfo.beginLoc, opInfo.varValue, /*structured=*/true,
+            /*implicit=*/false, opInfo.name, opInfo.bounds);
+        privateOp.setDataClause(mlir::acc::DataClause::acc_private);
+
+        {
+          mlir::OpBuilder::InsertionGuard guardCase(builder);
+          auto recipe = getOrCreateRecipe<mlir::acc::PrivateRecipeOp>(
+              cgf.getContext(), var, Decl::castToDeclContext(cgf.curFuncDecl),
+              opInfo.baseType, privateOp.getResult());
+          // TODO: OpenACC: The dialect is going to change in the near future to
+          // have these be on a different operation, so when that changes, we
+          // probably need to change these here.
+          operation.addPrivatization(builder.getContext(), privateOp, recipe);
+        }
+      }
+    } else if constexpr (isCombinedType<OpTy>) {
+      // Despite this being valid on ParallelOp or SerialOp, combined type
+      // applies to the 'loop'.
+      applyToLoopOp(clause);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitPrivateClause");
+    }
+  }
 };
 
 template <typename OpTy>
diff --git a/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
new file mode 100644
index 0000000000000..3306c55b4a8e7
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp
@@ -0,0 +1,522 @@
+// 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.private.recipe @privatization__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: } 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.private.recipe @privatization__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: 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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: } 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.private.recipe @privatization__ZTS14NonDefaultCtor : !cir.ptr<!rec_NonDefaultCtor> init {
+// CHECK-NEXT: ^bb0(%[[ARG:.*]]: !cir.ptr<!rec_NonDefaultCtor> {{.*}}):
+// CHECK-NEXT: cir.alloca !rec_NonDefaultCtor, !cir.ptr<!rec_NonDefaultCtor>, ["openacc.private.init"]
+// CHECK-NEXT: acc.yield
+// CHECK-NEXT: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+//
+// CHECK-NEXT: acc.private.recipe @privatization__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: }
+
+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 private(someInt)
+  for(int i = 0; i < 5; ++i);
+  // CHECK: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEINT]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "someInt"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTSi -> %[[PRIVATE]] : !cir.ptr<!s32i>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(someFloat)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[SOMEFLOAT]] : !cir.ptr<!cir.float>) -> !cir.ptr<!cir.float> {name = "someFloat"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTSf -> %[[PRIVATE]] : !cir.ptr<!cir.float>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+
+#pragma acc parallel loop private(noCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.parallel combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[NOCOPY]] : !cir.ptr<!rec_NoCopyConstruct>) -> !cir.ptr<!rec_NoCopyConstruct> {name = "noCopy"}
+  // CHECK-NEXT: acc.loop combined(parallel) private(@privatization__ZTS15NoCopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_NoCopyConstruct>
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEXT: } loc
+#pragma acc serial loop private(hasCopy)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: acc.serial combined(loop) {
+  // CHECK-NEXT: %[[PRIVATE:.*]] = acc.private varPtr(%[[HASCOPY]] : !cir.ptr<!rec_CopyConstruct>) -> !cir.ptr<!rec_CopyConstruct> {name = "hasCopy"}
+  // CHECK-NEXT: acc.loop combined(serial) private(@privatization__ZTS13CopyConstruct -> %[[PRIVATE]] : !cir.ptr<!rec_CopyConstruct>)
+  // CHECK: acc.yield
+  // CHECK-NEXT: } loc
+  // CHECK-NEXT: acc.yield
+  // CHECK-NEX...
[truncated]

erichkeane · 2025-07-30T16:17:23Z

@razvanlupusoru and @clementval : Would love review of the ACC dialect changes, but also confirmation that I got the IR correct. Particularly the lack of copy or constructor-initialization, and general privatization appertainment.

I realize this is going to change somewhat when Razvan's refactor patch gets merged (how is that going?), but wanted the feedback too.

To all:
The 3 tests are intentionally all-but-identical to ensure that this works the same/correct for each.

erichkeane · 2025-07-30T17:24:48Z

clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp

+};
+
+struct NonDefaultCtor {
+  NonDefaultCtor();


THESE examples in particular, it seems a little odd that this isn't necessary to call for privatization init, but I'm hoping @razvanlupusoru can explain that, or tell me I'm incorrect in my assumption.

It does indeed seem we had a slight misunderstanding. For a private variable we should be calling the default constructor - either the one that the user wrote or the implicit one compiler generated - and the call for this should reside in the acc.private.recipe's init region.

For firstprivate, it makes sense that it is OK just to enforce having a copy constructor and not necessarily having both default and copy constructors. However, the acc dialect uses another region to represent the copy - namely acc.firstprivate.recipe's copy region. So for a firstprivate, the body of the init region should look like you have it in the examples below (namely a simple cir.alloca) and the body of the copy region should have the call to the constructor.

However, this discussion makes me question if an acc dialect change would help here; namely, if it would be better if we had a single init region even for firstprivate instead of the two init and copy. I cannot think of any benefits of having the two. So I would welcome this change :)

Hi! Thanks! I definitely mis-understood a previous discussion. I'll work on getting the FE to do a default-init for the init section.

For firstprivate, it isn't clear why the copy is separate unless there is optimization opportunities, but having the FE put the copy operation in that region is easy enough.

andykaylor

A few nits, but otherwise looks good to me.

andykaylor · 2025-07-30T21:06:15Z

clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp

+                             DeclContext *dc, QualType baseType,
+                             mlir::Value mainOp) {
+    mlir::ModuleOp mod =
+        builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();


Suggested change

builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();

builder.getBlock()->getParentOfType<mlir::ModuleOp>();

getParentOfType walks the hierarchy until it finds the requested type.

mlir::Block doesn't seem to have getParentOfType? I think that must be why this uses getParent. OR am I missing something further?

clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp

razvanlupusoru · 2025-07-30T22:13:05Z

I realize this is going to change somewhat when Razvan's refactor patch gets merged (how is that going?), but wanted the feedback too.

I am hoping to finalize and submit this next week: #149210
Sorry about the delay. If your change is merged by then it is no problem - I should be able to resolve the conflicts because I am now building and running CIR tests.

erichkeane · 2025-07-31T17:10:29Z

I've been digging into the initialization part of this, and the patch ends up being pretty sizable. If it is OK with reviewers, I'd like to just submit this and do that as a separate review.

erichkeane · 2025-08-01T13:32:51Z

If I don't hear objections in the next few hours on this, I'll merge this and follow up with the patch to do 'init' properly.

That one ended up becoming somewhat large because it had a bit of work in Sema to do, so I'm keeping it separate to minimize the size of this one.

razvanlupusoru

Thank you!

Previously, llvm#151360 implemented 'private' clause lowering, but didn't properly initialize the variables. This patch adds that behavior to make sure we correctly get the constructor or other init called.

…1781) Previously, #151360 implemented 'private' clause lowering, but didn't properly initialize the variables. This patch adds that behavior to make sure we correctly get the constructor or other init called.

erichkeane requested review from razvanlupusoru and clementval July 30, 2025 16:15

erichkeane requested review from lanza, bcardosolopes, xlauko and andykaylor as code owners July 30, 2025 16:15

llvmbot added clang Clang issues not falling into any other category mlir mlir:openacc openacc ClangIR Anything related to the ClangIR project labels Jul 30, 2025

erichkeane commented Jul 30, 2025

View reviewed changes

andykaylor approved these changes Jul 30, 2025

View reviewed changes

Fix everything andy suggested

b40a8b4

razvanlupusoru approved these changes Aug 1, 2025

View reviewed changes

erichkeane merged commit 25c0776 into llvm:main Aug 1, 2025
9 checks passed

erichkeane mentioned this pull request Aug 1, 2025

[OpenACC][CIR] Implement 'init' lowering for private clause vars #151781

Merged

	builder.getBlock()->getParent()->getParentOfType<mlir::ModuleOp>();
	builder.getBlock()->getParentOfType<mlir::ModuleOp>();

[OpenACC][CIR] Implement 'private' clause lowering. #151360

[OpenACC][CIR] Implement 'private' clause lowering. #151360

Uh oh!

Conversation

erichkeane commented Jul 30, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jul 30, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

llvmbot commented Jul 30, 2025

Uh oh!

llvmbot commented Jul 30, 2025

Uh oh!

llvmbot commented Jul 30, 2025

Uh oh!

erichkeane commented Jul 30, 2025

Uh oh!

erichkeane Jul 30, 2025

Choose a reason for hiding this comment

Uh oh!

razvanlupusoru Jul 30, 2025

Choose a reason for hiding this comment

Uh oh!

erichkeane Jul 31, 2025

Choose a reason for hiding this comment

Uh oh!

andykaylor left a comment

Choose a reason for hiding this comment

Uh oh!

andykaylor Jul 30, 2025

Choose a reason for hiding this comment

Uh oh!

erichkeane Jul 31, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

razvanlupusoru commented Jul 30, 2025

Uh oh!

erichkeane commented Jul 31, 2025

Uh oh!

erichkeane commented Aug 1, 2025

Uh oh!

razvanlupusoru left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

erichkeane commented Jul 30, 2025 •

edited

Loading

llvmbot commented Jul 30, 2025 •

edited

Loading