[clang] [CIR] Add cir.builtin_int_cast operation to replace uses of builtin.unrealized_conversion_cast (PR #201592)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Jun 4 07:03:12 PDT 2026
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
More information about the cfe-commits
mailing list