[Mlir-commits] [clang] [mlir] [OpenACC][CIR] Implement 'private' clause lowering. (PR #151360)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Wed Jul 30 09:16:00 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-mlir-openacc
Author: Erich Keane (erichkeane)
<details>
<summary>Changes</summary>
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
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]
``````````
</details>
https://github.com/llvm/llvm-project/pull/151360
More information about the Mlir-commits
mailing list