llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clangir Author: Jan Leyonberg (jsjodin) <details> <summary>Changes</summary> This patch adds a new operation builtin_int_cast to handle casting between CIR integer types and builtin integer types. This will replace the current use of the builtin.unrealized_conversion_cast since this operation is only intended to be used temporarily when doing transformations. Assisted-by: Cursor/Claude Opus 4.8 High --- Patch is 1.07 MiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/201592.diff 74 Files Affected: - (modified) clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h (+9) - (modified) clang/include/clang/CIR/Dialect/IR/CIROps.td (+38) - (modified) clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp (+1-3) - (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp (+9-13) - (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+1-3) - (modified) clang/lib/CIR/Dialect/IR/CIRDialect.cpp (+46) - (modified) clang/lib/CIR/Dialect/Transforms/CIRCanonicalize.cpp (+7-7) - (modified) clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp (+25) - (modified) clang/test/CIR/CodeGenOpenACC/cache.c (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/combined-copy.c (+79-79) - (modified) clang/test/CIR/CodeGenOpenACC/combined-copy.cpp (+16-16) - (modified) clang/test/CIR/CodeGenOpenACC/combined-firstprivate-clause.cpp (+50-50) - (modified) clang/test/CIR/CodeGenOpenACC/combined-private-clause.cpp (+40-40) - (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-default-ops.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-float.cpp (+24-24) - (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-inline-ops.cpp (+54-54) - (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-int.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/combined-reduction-clause-outline-ops.cpp (+54-54) - (modified) clang/test/CIR/CodeGenOpenACC/combined.cpp (+91-91) - (modified) clang/test/CIR/CodeGenOpenACC/compute-copy.c (+79-79) - (modified) clang/test/CIR/CodeGenOpenACC/compute-copy.cpp (+16-16) - (modified) clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.c (+24-24) - (modified) clang/test/CIR/CodeGenOpenACC/compute-firstprivate-clause.cpp (+50-50) - (modified) clang/test/CIR/CodeGenOpenACC/compute-private-clause.c (+18-18) - (modified) clang/test/CIR/CodeGenOpenACC/compute-private-clause.cpp (+40-40) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-default-ops.c (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-default-ops.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-float.c (+24-24) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-float.cpp (+24-24) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-inline-ops.cpp (+54-54) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-int.c (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-int.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-outline-ops.cpp (+54-54) - (modified) clang/test/CIR/CodeGenOpenACC/compute-reduction-clause-unsigned-int.c (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/data.c (+28-28) - (modified) clang/test/CIR/CodeGenOpenACC/declare-copy.cpp (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/declare-copyin.cpp (+28-28) - (modified) clang/test/CIR/CodeGenOpenACC/declare-copyout.cpp (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/declare-create.cpp (+28-28) - (modified) clang/test/CIR/CodeGenOpenACC/declare-deviceresident.cpp (+28-28) - (modified) clang/test/CIR/CodeGenOpenACC/declare-link.cpp (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/declare-present.cpp (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/enter-data.c (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/exit-data.c (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/firstprivate-clause-recipes.cpp (+42-42) - (modified) clang/test/CIR/CodeGenOpenACC/host_data.c (+2-2) - (modified) clang/test/CIR/CodeGenOpenACC/init.c (+3-3) - (modified) clang/test/CIR/CodeGenOpenACC/kernels.c (+58-58) - (modified) clang/test/CIR/CodeGenOpenACC/loop-private-clause.cpp (+40-40) - (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-default-ops.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-float.cpp (+24-24) - (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-inline-ops.cpp (+54-54) - (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-int.cpp (+36-36) - (modified) clang/test/CIR/CodeGenOpenACC/loop-reduction-clause-outline-ops.cpp (+54-54) - (modified) clang/test/CIR/CodeGenOpenACC/loop.cpp (+22-22) - (modified) clang/test/CIR/CodeGenOpenACC/parallel.c (+69-69) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-CtorDtor.cpp (+32-32) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-array-recipes-NoOps.cpp (+16-16) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-array-recipes-CtorDtor.cpp (+108-108) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-array-recipes-NoOps.cpp (+66-66) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-array-recipes-int.cpp (+24-24) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-recipes-CtorDtor.cpp (+34-34) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-recipes-NoOps.cpp (+22-22) - (modified) clang/test/CIR/CodeGenOpenACC/private-clause-pointer-recipes-int.cpp (+10-10) - (modified) clang/test/CIR/CodeGenOpenACC/reduction-clause-recipes.cpp (+60-60) - (modified) clang/test/CIR/CodeGenOpenACC/serial.c (+28-28) - (modified) clang/test/CIR/CodeGenOpenACC/set.c (+5-5) - (modified) clang/test/CIR/CodeGenOpenACC/shutdown.c (+3-3) - (modified) clang/test/CIR/CodeGenOpenACC/update.c (+12-12) - (modified) clang/test/CIR/CodeGenOpenACC/wait.c (+14-14) - (added) clang/test/CIR/IR/builtin-int-cast.cir (+50) - (added) clang/test/CIR/IR/invalid-builtin-int-cast.cir (+37) - (added) clang/test/CIR/Lowering/builtin-int-cast.cir (+56) - (added) clang/test/CIR/Transforms/builtin-int-cast-fold.cir (+29) ``````````diff diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index a29fb45e95032..681e129e41cb7 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -529,6 +529,15 @@ class CIRBaseBuilderTy : public mlir::OpBuilder { return createCast(cir::CastKind::integral, src, newTy); } + mlir::Value createBuiltinIntCast(mlir::Location loc, mlir::Value src, + mlir::Type newTy) { + return cir::BuiltinIntCastOp::create(*this, loc, newTy, src); + } + + mlir::Value createBuiltinIntCast(mlir::Value src, mlir::Type newTy) { + return createBuiltinIntCast(src.getLoc(), src, newTy); + } + mlir::Value createIntToPtr(mlir::Value src, mlir::Type newTy) { return createCast(cir::CastKind::int_to_ptr, src, newTy); } diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index c4d08d5337031..0dfd416bf24dd 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -288,6 +288,44 @@ def CIR_CastOp : CIR_Op<"cast", [ } +//===----------------------------------------------------------------------===// +// BuiltinIntCastOp +//===----------------------------------------------------------------------===// + +def CIR_IntOrBuiltinIntType : AnyTypeOf<[CIR_IntType, AnyInteger, Index]>; + +def CIR_BuiltinIntCastOp : CIR_Op<"builtin_int_cast", [Pure]> { + let summary = "Cast between a CIR integer and a builtin integer"; + let description = [{ + Convert between a CIR integer type (`!cir.int`) and a builtin MLIR integer + type (`AnyInteger`, e.g. `i32`, `si32`, `ui32`) or `index`, and vice versa. + + This allows using operations from e.g. OpenMP or OpenACC dialects + that expect the builtin types with CIR operations. Casting can be done + in either direction. + + Example: + + ```mlir + // CIR integer cast to a builtin integer. + %0 = cir.builtin_int_cast %ciri : !cir.int<s, 32> -> i32 + + // Builtin induction variable / bound cast to CIR type. + %1 = cir.builtin_int_cast %iv : index -> !cir.int<u, 64> + ``` + }]; + + let arguments = (ins CIR_IntOrBuiltinIntType:$src); + let results = (outs CIR_IntOrBuiltinIntType:$result); + + let assemblyFormat = [{ + $src `:` type($src) `->` type($result) attr-dict + }]; + + let hasVerifier = 1; + let hasFolder = 1; +} + //===----------------------------------------------------------------------===// // DynamicCastOp //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp index e7bf3bcc85c0b..e9725940dd501 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACC.cpp @@ -48,9 +48,7 @@ mlir::Value CIRGenFunction::emitOpenACCIntExpr(const Expr *intExpr) { ? mlir::IntegerType::SignednessSemantics::Signed : mlir::IntegerType::SignednessSemantics::Unsigned); - auto conversionOp = mlir::UnrealizedConversionCastOp::create( - builder, exprLoc, targetType, expr); - return conversionOp.getResult(0); + return builder.createBuiltinIntCast(exprLoc, expr, targetType); } mlir::Value CIRGenFunction::createOpenACCConstantInt(mlir::Location loc, diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp index eab045e186699..7018fdd995691 100644 --- a/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCRecipe.cpp @@ -126,9 +126,7 @@ mlir::Value OpenACCRecipeBuilderBase::makeBoundsAlloca( auto getUpperBound = [&](mlir::Value bound) { auto upperBoundVal = mlir::acc::GetUpperboundOp::create(builder, loc, idxType, bound); - return mlir::UnrealizedConversionCastOp::create(builder, loc, itrTy, - upperBoundVal.getResult()) - .getResult(0); + return builder.createBuiltinIntCast(loc, upperBoundVal.getResult(), itrTy); }; auto isArrayTy = [&](QualType ty) { @@ -253,12 +251,12 @@ std::pair<mlir::Value, mlir::Value> OpenACCRecipeBuilderBase::createBoundsLoop( // get the lower and upper bound for iterating over. auto lowerBoundVal = mlir::acc::GetLowerboundOp::create(builder, loc, idxType, bound); - auto lbConversion = mlir::UnrealizedConversionCastOp::create( - builder, loc, itrTy, lowerBoundVal.getResult()); + mlir::Value lbConversion = + builder.createBuiltinIntCast(loc, lowerBoundVal.getResult(), itrTy); auto upperBoundVal = mlir::acc::GetUpperboundOp::create(builder, loc, idxType, bound); - auto ubConversion = mlir::UnrealizedConversionCastOp::create( - builder, loc, itrTy, upperBoundVal.getResult()); + mlir::Value ubConversion = + builder.createBuiltinIntCast(loc, upperBoundVal.getResult(), itrTy); // Create a memory location for the iterator. auto itr = @@ -268,20 +266,18 @@ std::pair<mlir::Value, mlir::Value> OpenACCRecipeBuilderBase::createBoundsLoop( if (inverse) { cir::ConstantOp constOne = builder.getConstInt(loc, itrTy, 1); - auto sub = - cir::SubOp::create(builder, loc, ubConversion.getResult(0), constOne); + auto sub = cir::SubOp::create(builder, loc, ubConversion, constOne); // Upperbound is exclusive, so subtract 1. builder.CIRBaseBuilderTy::createStore(loc, sub, itr); } else { // Lowerbound is inclusive, so we can include it. - builder.CIRBaseBuilderTy::createStore(loc, lbConversion.getResult(0), - itr); + builder.CIRBaseBuilderTy::createStore(loc, lbConversion, itr); } // Save the 'end' iterator based on whether we are inverted or not. This // end iterator never changes, so we can just get it and convert it, so no // need to store/load/etc. - auto endItr = inverse ? lbConversion : ubConversion; + mlir::Value endItr = inverse ? lbConversion : ubConversion; builder.createFor( loc, @@ -291,7 +287,7 @@ std::pair<mlir::Value, mlir::Value> OpenACCRecipeBuilderBase::createBoundsLoop( // Use 'not equal' since we are just doing an increment/decrement. auto cmp = builder.createCompare( loc, inverse ? cir::CmpOpKind::ge : cir::CmpOpKind::lt, loadCur, - endItr.getResult(0)); + endItr); builder.createCondition(cmp); }, /*bodyBuilder=*/ diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp index 11aad17187fbc..7f59cccec2eb9 100644 --- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp @@ -187,9 +187,7 @@ CIRGenFunction::emitOpenACCWaitConstruct(const OpenACCWaitConstruct &s) { ? mlir::IntegerType::SignednessSemantics::Signed : mlir::IntegerType::SignednessSemantics::Unsigned); - auto conversionOp = mlir::UnrealizedConversionCastOp::create( - builder, exprLoc, targetType, expr); - return conversionOp.getResult(0); + return builder.createBuiltinIntCast(exprLoc, expr, targetType); }; // Emit the correct 'wait' clauses. diff --git a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp index cf07fc4f0833a..61db79b5c7bb0 100644 --- a/clang/lib/CIR/Dialect/IR/CIRDialect.cpp +++ b/clang/lib/CIR/Dialect/IR/CIRDialect.cpp @@ -17,6 +17,7 @@ #include "clang/CIR/Dialect/IR/CIRTypes.h" #include "mlir/IR/Attributes.h" +#include "mlir/IR/BuiltinTypes.h" #include "mlir/IR/DialectImplementation.h" #include "mlir/IR/PatternMatch.h" #include "mlir/IR/Value.h" @@ -972,6 +973,51 @@ OpFoldResult cir::CastOp::fold(FoldAdaptor adaptor) { return {}; } +//===----------------------------------------------------------------------===// +// BuiltinIntCastOp +//===----------------------------------------------------------------------===// + +LogicalResult cir::BuiltinIntCastOp::verify() { + mlir::Type srcType = getSrc().getType(); + mlir::Type resType = getType(); + + auto srcCirInt = mlir::dyn_cast<cir::IntType>(srcType); + auto resCirInt = mlir::dyn_cast<cir::IntType>(resType); + + // One side must be a CIR integer the other must be a builtin + // integer or index type. + if (static_cast<bool>(srcCirInt) == static_cast<bool>(resCirInt)) + return emitOpError() + << "requires exactly one '!cir.int' operand or result; the other " + "must be a builtin integer or 'index' type"; + + mlir::Type builtinType = srcCirInt ? resType : srcType; + if (!mlir::isa<mlir::IntegerType, mlir::IndexType>(builtinType)) + return emitOpError() << "requires a builtin integer or 'index' type on the " + "non-CIR side"; + + // The cast preserves bit width. 'index' has no fixed width, so only check + // when the builtin side is a fixed-width integer. + if (auto builtinInt = mlir::dyn_cast<mlir::IntegerType>(builtinType)) { + cir::IntType cirInt = srcCirInt ? srcCirInt : resCirInt; + if (cirInt.getWidth() != builtinInt.getWidth()) + return emitOpError() + << "requires the CIR and builtin integer types to have the same " + "width; use 'cir.cast' for width conversions"; + } + + return success(); +} + +OpFoldResult cir::BuiltinIntCastOp::fold(FoldAdaptor adaptor) { + // Fold: builtin_int_cast(builtin_int_cast(x)) -> x + // Inner source type must match the cast's result type. + if (auto inner = getSrc().getDefiningOp<cir::BuiltinIntCastOp>()) + if (inner.getSrc().getType() == getType()) + return inner.getSrc(); + return {}; +} + //===----------------------------------------------------------------------===// // CallOp //===----------------------------------------------------------------------===// diff --git a/clang/lib/CIR/Dialect/Transforms/CIRCanonicalize.cpp b/clang/lib/CIR/Dialect/Transforms/CIRCanonicalize.cpp index da08f21977066..7914f0068f720 100644 --- a/clang/lib/CIR/Dialect/Transforms/CIRCanonicalize.cpp +++ b/clang/lib/CIR/Dialect/Transforms/CIRCanonicalize.cpp @@ -70,13 +70,13 @@ void CIRCanonicalizePass::runOnOperation() { // Many operations are here to perform a manual `fold` in // applyOpPatternsGreedily. - if (isa<BrOp, BrCondOp, CastOp, ScopeOp, SwitchOp, SelectOp, IncOp, DecOp, - MinusOp, NotOp, AddOp, MulOp, AndOp, OrOp, XorOp, MaxOp, MinOp, - ComplexCreateOp, ComplexImagOp, ComplexRealOp, VecCmpOp, - VecCreateOp, VecExtractOp, VecShuffleOp, VecShuffleDynamicOp, - VecTernaryOp, BitClrsbOp, BitClzOp, BitCtzOp, BitFfsOp, BitParityOp, - BitPopcountOp, BitReverseOp, ByteSwapOp, RotateOp, ConstantOp, - CleanupScopeOp>(op)) + if (isa<BrOp, BrCondOp, CastOp, BuiltinIntCastOp, ScopeOp, SwitchOp, + SelectOp, IncOp, DecOp, MinusOp, NotOp, AddOp, MulOp, AndOp, OrOp, + XorOp, MaxOp, MinOp, ComplexCreateOp, ComplexImagOp, ComplexRealOp, + VecCmpOp, VecCreateOp, VecExtractOp, VecShuffleOp, + VecShuffleDynamicOp, VecTernaryOp, BitClrsbOp, BitClzOp, BitCtzOp, + BitFfsOp, BitParityOp, BitPopcountOp, BitReverseOp, ByteSwapOp, + RotateOp, ConstantOp, CleanupScopeOp>(op)) ops.push_back(op); }); diff --git a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp index 8c7e1406d6567..1a717f2af5888 100644 --- a/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp +++ b/clang/lib/CIR/Lowering/DirectToLLVM/LowerToLLVM.cpp @@ -1437,6 +1437,31 @@ mlir::LogicalResult CIRToLLVMCastOpLowering::matchAndRewrite( return mlir::success(); } +mlir::LogicalResult CIRToLLVMBuiltinIntCastOpLowering::matchAndRewrite( + cir::BuiltinIntCastOp op, OpAdaptor adaptor, + mlir::ConversionPatternRewriter &rewriter) const { + // Both the CIR integer and the builtin integer/index lower to LLVM integer + // types, so this cast becomes an integer resize. Signedness is taken from + // the CIR integer side (the builtin/index side is treated as signless). + bool isUnsigned = true; + if (auto cirSrc = mlir::dyn_cast<cir::IntType>(op.getSrc().getType())) + isUnsigned = cirSrc.isUnsigned(); + else if (auto cirDst = mlir::dyn_cast<cir::IntType>(op.getType())) + isUnsigned = cirDst.isUnsigned(); + + mlir::Value llvmSrc = adaptor.getSrc(); + mlir::Type llvmDstTy = getTypeConverter()->convertType(op.getType()); + auto srcIntTy = mlir::cast<mlir::IntegerType>(llvmSrc.getType()); + auto dstIntTy = mlir::cast<mlir::IntegerType>(llvmDstTy); + + // For equal widths getLLVMIntCast returns the source unchanged, which makes + // the common same-width cast a no-op. + rewriter.replaceOp(op, + getLLVMIntCast(rewriter, llvmSrc, dstIntTy, isUnsigned, + srcIntTy.getWidth(), dstIntTy.getWidth())); + return mlir::success(); +} + static mlir::Value convertToIndexTy(mlir::ConversionPatternRewriter &rewriter, mlir::ModuleOp mod, mlir::Value index, mlir::Type baseTy, cir::IntType strideTy) { diff --git a/clang/test/CIR/CodeGenOpenACC/cache.c b/clang/test/CIR/CodeGenOpenACC/cache.c index d82230a5e3841..76662782d6daa 100644 --- a/clang/test/CIR/CodeGenOpenACC/cache.c +++ b/clang/test/CIR/CodeGenOpenACC/cache.c @@ -27,7 +27,7 @@ void acc_cache() { } } // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 @@ -35,9 +35,9 @@ void acc_cache() { // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false} // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5> - // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si32 + // CHECK-NEXT: %[[FIVE_CAST:.*]] = cir.builtin_int_cast %[[FIVE]] : !s32i -> si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) @@ -52,7 +52,7 @@ void acc_cache() { #pragma acc cache(iArr[1], fArr[1:5]) } // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 @@ -60,9 +60,9 @@ void acc_cache() { // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false} // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5> - // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si32 + // CHECK-NEXT: %[[FIVE_CAST:.*]] = cir.builtin_int_cast %[[FIVE]] : !s32i -> si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) @@ -78,7 +78,7 @@ void acc_cache() { } // CHECK-NEXT: acc.parallel combined(loop) { // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 @@ -86,9 +86,9 @@ void acc_cache() { // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false} // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5> - // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si32 + // CHECK-NEXT: %[[FIVE_CAST:.*]] = cir.builtin_int_cast %[[FIVE]] : !s32i -> si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) @@ -108,7 +108,7 @@ void acc_cache() { } // CHECK-NEXT: acc.parallel combined(loop) { // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST2:.*]] = arith.constant 1 : i64 @@ -116,9 +116,9 @@ void acc_cache() { // CHECK-NEXT: %[[CACHE1:.*]] = acc.cache varPtr(%[[IARR]] : !cir.ptr<!cir.array<!s32i x 10>>) bounds(%[[BOUNDS]]) -> !cir.ptr<!cir.array<!s32i x 10>> {name = "iArr[1]", structured = false} // // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[FIVE:.*]] = cir.const #cir.int<5> - // CHECK-NEXT: %[[FIVE_CAST:.*]] = builtin.unrealized_conversion_cast %[[FIVE]] : !s32i to si32 + // CHECK-NEXT: %[[FIVE_CAST:.*]] = cir.builtin_int_cast %[[FIVE]] : !s32i -> si32 // CHECK-NEXT: %[[ZERO_CONST:.*]] = arith.constant 0 : i64 // CHECK-NEXT: %[[ONE_CONST:.*]] = arith.constant 1 : i64 // CHECK-NEXT: %[[BOUNDS:.*]] = acc.bounds lowerbound(%[[ONE_CAST]] : si32) extent(%[[FIVE_CAST]] : si32) stride(%[[ONE_CONST]] : i64) startIdx(%[[ZERO_CONST]] : i64) diff --git a/clang/test/CIR/CodeGenOpenACC/combined-copy.c b/clang/test/CIR/CodeGenOpenACC/combined-copy.c index f16c25188bd1b..f03eb8ff7dd65 100644 --- a/clang/test/CIR/CodeGenOpenACC/combined-copy.c +++ b/clang/test/CIR/CodeGenOpenACC/combined-copy.c @@ -145,7 +145,7 @@ void acc_compute(int parmVar) { #pragma acc kernels loop copy(localVar1, localVar2) async(1) for(int i = 0; i < 5; ++i); // CHECK-NEXT: %[[ONE:.*]] = cir.const #cir.int<1> : !s32i - // CHECK-NEXT: %[[ONE_CAST:.*]] = builtin.unrealized_conversion_cast %[[ONE]] : !s32i to si32 + // CHECK-NEXT: %[[ONE_CAST:.*]] = cir.builtin_int_cast %[[ONE]] : !s32i -> si32 // CHECK-NEXT: %[[COPYIN1:.*]] = acc.copyin varPtr(%[[LOCAL1]] : !cir.ptr<!s32i>) async(%[[ONE_CAST]] : si32) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_copy>, name = "localVar1"} loc // CH... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/201592 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
