diff --git a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp index 091441d564334..a8b7f60b6a014 100644 --- a/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp +++ b/clang/lib/CIR/Dialect/Transforms/TargetLowering/CIRABIRewriteContext.cpp @@ -282,17 +282,16 @@ mlir::ArrayAttr updateResAttrs(mlir::MLIRContext *ctx, /// is written through a source-typed view and returned as a destination-typed /// view. /// -/// The temporary alloca is placed at the start of the enclosing function's -/// entry block so that it composes correctly with the HoistAllocas pass -/// regardless of pipeline ordering. +/// The temporary alloca is placed at the start of \p slotBlock, which must +/// dominate every use of the coerced value and must be a block that ends up +/// inside the enclosing function's entry block after any later outlining. /// /// Any operations the helper creates are appended to \p createdOps so the /// caller can pass them to replaceAllUsesExcept and avoid clobbering the /// store's value operand when later rewiring the source value. mlir::Value emitCoercionToMemory(mlir::OpBuilder &builder, mlir::Location loc, - mlir::Type dstTy, mlir::Value src, - mlir::FunctionOpInterface funcOp, + mlir::Type dstTy, mlir::Value src, mlir::Block *slotBlock, const mlir::DataLayout &dl, SmallPtrSetImpl &createdOps) { mlir::Type srcTy = src.getType(); @@ -312,8 +311,7 @@ emitCoercionToMemory(mlir::OpBuilder &builder, mlir::Location loc, cir::AllocaOp alloca; { mlir::OpBuilder::InsertionGuard guard(builder); - mlir::Block &entry = funcOp->getRegion(0).front(); - builder.setInsertionPointToStart(&entry); + builder.setInsertionPointToStart(slotBlock); alloca = cir::AllocaOp::create(builder, loc, slotPtrTy, builder.getStringAttr("coerce"), builder.getI64IntegerAttr(allocaAlign)); @@ -346,11 +344,10 @@ emitCoercionToMemory(mlir::OpBuilder &builder, mlir::Location loc, /// load of the destination-typed view. mlir::Value emitCoercion(mlir::OpBuilder &builder, mlir::Location loc, mlir::Type dstTy, mlir::Value src, - mlir::FunctionOpInterface funcOp, - const mlir::DataLayout &dl, + mlir::Block *slotBlock, const mlir::DataLayout &dl, SmallPtrSetImpl &createdOps) { mlir::Value dstSlot = - emitCoercionToMemory(builder, loc, dstTy, src, funcOp, dl, createdOps); + emitCoercionToMemory(builder, loc, dstTy, src, slotBlock, dl, createdOps); auto load = cir::LoadOp::create(builder, loc, dstSlot); createdOps.insert(load); return load; @@ -360,10 +357,31 @@ mlir::Value emitCoercion(mlir::OpBuilder &builder, mlir::Location loc, /// (e.g. call-site coercion where we don't replaceAllUsesExcept). mlir::Value emitCoercion(mlir::OpBuilder &builder, mlir::Location loc, mlir::Type dstTy, mlir::Value src, - mlir::FunctionOpInterface funcOp, - const mlir::DataLayout &dl) { + mlir::Block *slotBlock, const mlir::DataLayout &dl) { SmallPtrSet ignored; - return emitCoercion(builder, loc, dstTy, src, funcOp, dl, ignored); + return emitCoercion(builder, loc, dstTy, src, slotBlock, dl, ignored); +} + +/// The block a coercion slot's alloca belongs at the start of. +/// +/// Normally the enclosing function's entry block, where HoistAllocas expects +/// allocas to be. A body carrying a call is not always inside a function +/// when this pass runs, though, because LoweringPrepare runs after it: a +/// namespace-scope `T g = makeT();` is still in its cir.global ctor region, +/// and an OpenACC recipe's init and destroy bodies are in regions the module +/// owns. Those take the outermost region below the module, which dominates +/// the whole body and travels with it when the body is outlined. +mlir::Block *coercionSlotBlock(mlir::Operation *op) { + if (auto funcOp = op->getParentOfType()) + return &funcOp->getRegion(0).front(); + mlir::Region *region = op->getParentRegion(); + while (mlir::Region *outer = region->getParentRegion()) { + if (mlir::isa(outer->getParentOp())) + break; + region = outer; + } + assert(!region->empty() && "coercion slot needs a block to hold the alloca"); + return ®ion->front(); } /// Insert coercion before each cir.return so the returned value matches the @@ -382,7 +400,8 @@ void insertReturnCoercion(mlir::FunctionOpInterface funcOp, continue; builder.setInsertionPoint(r); mlir::Value coerced = - emitCoercion(builder, r.getLoc(), coercedRetTy, origVal, funcOp, dl); + emitCoercion(builder, r.getLoc(), coercedRetTy, origVal, + &funcOp->getRegion(0).front(), dl); r->setOperand(0, coerced); } } @@ -579,7 +598,7 @@ void insertArgCoercion(mlir::FunctionOpInterface funcOp, Value finalVal = flatLoaded; if (origTy != flatTy) { SmallPtrSet coercionOps; - finalVal = emitCoercion(builder, loc, origTy, flatLoaded, funcOp, dl, + finalVal = emitCoercion(builder, loc, origTy, flatLoaded, &entry, dl, coercionOps); flattenOps.insert(coercionOps.begin(), coercionOps.end()); } @@ -604,7 +623,7 @@ void insertArgCoercion(mlir::FunctionOpInterface funcOp, builder.setInsertionPointToStart(&entry); SmallPtrSet coercionOps; mlir::Value adapted = emitCoercion(builder, funcOp.getLoc(), oldArgTy, - blockArg, funcOp, dl, coercionOps); + blockArg, &entry, dl, coercionOps); // Replace blockArg uses with the adapted value, except inside the // helper ops we just created. This is critical: the StoreOp's value @@ -1122,7 +1141,7 @@ CIRABIRewriteContext::rewriteCallSite(mlir::Operation *callOp, auto call = mlir::cast(callOp); mlir::MLIRContext *ctx = callOp->getContext(); - auto enclosingFunc = call->getParentOfType(); + mlir::Block *slotBlock = coercionSlotBlock(call); builder.setInsertionPoint(call); @@ -1153,9 +1172,8 @@ CIRABIRewriteContext::rewriteCallSite(mlir::Operation *callOp, // alloca when possible). if (arg.getType() != flatTy) { SmallPtrSet coercionOps; - mlir::Value coercedPtr = - emitCoercionToMemory(builder, call.getLoc(), flatTy, arg, - enclosingFunc, dl, coercionOps); + mlir::Value coercedPtr = emitCoercionToMemory( + builder, call.getLoc(), flatTy, arg, slotBlock, dl, coercionOps); for (auto [f, fieldTy] : llvm::enumerate(flatTy.getMembers())) { mlir::Type fieldPtrTy = cir::PointerType::get(fieldTy); auto fieldPtr = @@ -1178,8 +1196,8 @@ CIRABIRewriteContext::rewriteCallSite(mlir::Operation *callOp, replacedWholeLoads); } else if (ac.kind == ArgKind::Direct && ac.coercedType && arg.getType() != ac.coercedType) { - arg = emitCoercion(builder, call.getLoc(), ac.coercedType, arg, - enclosingFunc, dl); + arg = emitCoercion(builder, call.getLoc(), ac.coercedType, arg, slotBlock, + dl); newArgs.push_back(arg); } else if (ac.kind == ArgKind::Indirect) { // byval and byref: allocate a stack slot, copy the value in, and pass @@ -1237,9 +1255,8 @@ CIRABIRewriteContext::rewriteCallSite(mlir::Operation *callOp, // emit a coercion back to the original type for the call's existing uses. if (returnNeedsCoercion) { builder.setInsertionPointAfter(newCall); - mlir::Value coercedBack = - emitCoercion(builder, call.getLoc(), origRetTy, newCall.getResult(), - enclosingFunc, dl); + mlir::Value coercedBack = emitCoercion(builder, call.getLoc(), origRetTy, + newCall.getResult(), slotBlock, dl); call.getResult().replaceAllUsesWith(coercedBack); } diff --git a/clang/test/CIR/CodeGen/global-init-coerced-return.cpp b/clang/test/CIR/CodeGen/global-init-coerced-return.cpp new file mode 100644 index 0000000000000..a7b4b3abfb6a2 --- /dev/null +++ b/clang/test/CIR/CodeGen/global-init-coerced-return.cpp @@ -0,0 +1,32 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -clangir-enable-call-conv-lowering -emit-cir %s -o %t.cir +// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -fclangir -clangir-enable-call-conv-lowering -emit-llvm %s -o %t-cir.ll +// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm %s -o %t.ll +// RUN: FileCheck --check-prefix=OGCG --input-file=%t.ll %s + +struct Pair { long a, b; }; +Pair makePair(); + +// The initializer is still in the cir.global ctor region when CallConvLowering +// runs, so its coercion slot has no enclosing cir.func to be placed in. +Pair g = makePair(); + +// CIR: cir.func private @_Z8makePairv() -> !rec_anon_struct + +// The slot lands first in the block the initializer is later outlined into. +// CIR-LABEL: cir.func internal private @__cxx_global_var_init() +// CIR-NEXT: %[[COERCE:.+]] = cir.alloca "coerce" align(8) : !cir.ptr +// CIR: %[[RET:.+]] = cir.call @_Z8makePairv() : () -> !rec_anon_struct +// CIR-NEXT: cir.store %[[RET]], %[[COERCE]] : !rec_anon_struct, !cir.ptr + +// LLVM-LABEL: define internal void @__cxx_global_var_init() +// LLVM-NEXT: %[[COERCE:.+]] = alloca { i64, i64 }, i64 1, align 8 +// LLVM-NEXT: %[[RET:.+]] = call { i64, i64 } @_Z8makePairv() +// LLVM-NEXT: store { i64, i64 } %[[RET]], ptr %[[COERCE]], align 8 + +// Both backends take the register pair back. Classic pulls the eightbytes out +// with extractvalue where CIR reads them back through the slot. +// OGCG-LABEL: define internal void @__cxx_global_var_init() +// OGCG: %[[RET:.+]] = call { i64, i64 } @_Z8makePairv() +// OGCG-NEXT: %{{.+}} = extractvalue { i64, i64 } %[[RET]], 0 diff --git a/clang/test/CIR/CodeGenOpenACC/private-clause-coerced-ctor.cpp b/clang/test/CIR/CodeGenOpenACC/private-clause-coerced-ctor.cpp new file mode 100644 index 0000000000000..9652c6ce55e2f --- /dev/null +++ b/clang/test/CIR/CodeGenOpenACC/private-clause-coerced-ctor.cpp @@ -0,0 +1,30 @@ +// RUN: %clang_cc1 -fopenacc -triple x86_64-unknown-linux-gnu -fclangir \ +// RUN: -clangir-enable-call-conv-lowering -emit-cir %s -o - | FileCheck %s + +struct Pair { + long a, b; +}; + +// Pair is two eightbytes, so the constructor's argument is coerced and the +// call needs a coercion slot. The recipe body is not inside a cir.func when +// CallConvLowering runs, so the slot goes at the start of the init region. +struct HasCoercedCtor { + HasCoercedCtor(Pair p = Pair{}); + ~HasCoercedCtor(); +}; + +void privatized() { + HasCoercedCtor c; +#pragma acc parallel private(c) + ; +} + +// CHECK: acc.private.recipe @privatization__ZTS14HasCoercedCtor : !cir.ptr init { +// CHECK-NEXT: ^bb0(%{{[^,)]+}}: !cir.ptr +// CHECK-NEXT: %[[COERCE:.*]] = cir.alloca "coerce" {{.*}} : !cir.ptr +// CHECK-NEXT: cir.alloca "openacc.private.init" +// CHECK: cir.store %{{.+}}, %[[COERCE]] : !rec_Pair, !cir.ptr +// CHECK: %[[VIEW:.*]] = cir.cast bitcast %[[COERCE]] : !cir.ptr -> !cir.ptr +// CHECK: %[[A:.*]] = cir.load %{{.+}} : !cir.ptr, !s64i +// CHECK: %[[B:.*]] = cir.load %{{.+}} : !cir.ptr, !s64i +// CHECK: cir.call @_ZN14HasCoercedCtorC1E4Pair(%{{.+}}, %[[A]], %[[B]]) : (!cir.ptr {{.*}}, !s64i, !s64i) -> ()