Skip to content

[CIR] Add cir.builtin_int_cast operation to replace uses of builtin.unrealized_conversion_cast#201592

Open
jsjodin wants to merge 3 commits into
llvm:mainfrom
jsjodin:jleyonberg/cir-integer-cast
Open

[CIR] Add cir.builtin_int_cast operation to replace uses of builtin.unrealized_conversion_cast#201592
jsjodin wants to merge 3 commits into
llvm:mainfrom
jsjodin:jleyonberg/cir-integer-cast

Conversation

@jsjodin
Copy link
Copy Markdown
Contributor

@jsjodin jsjodin commented Jun 4, 2026

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

jsjodin added 3 commits June 4, 2026 09:39
@jsjodin
[CIR] Add cir.builtin_int_cast operation
705de6f 
This patch adds a new operation to allow casting between CIR integer types
and builtin integer/index types. This operation works in both directions.
@jsjodin
[CIR] Add lowering of cli.builtin_int_cast to LLVM IR dialect
8199ff0 
This patch adds lowering of the cli.builtin_int_cast to the LLVM IR dialect.
@jsjodin
[CIR] Change OpenACC codegen to use the cir.builtin_int_cast op
c891a14 
This patch changes the OpenACC codegen to use the cir.builtin_int_cast op
instead of the unrealized conversion casts.
@jsjodin jsjodin requested a review from erichkeane June 4, 2026 14:02
@llvmorg-github-actions llvmorg-github-actions Bot added clang Clang issues not falling into any other category ClangIR Anything related to the ClangIR project labels Jun 4, 2026
@llvmorg-github-actions
Copy link
Copy Markdown

llvmorg-github-actions Bot commented Jun 4, 2026 •
edited
Loading

@llvm/pr-subscribers-clang

@llvm/pr-subscribers-clangir

Author: Jan Leyonberg (jsjodin)

Changes

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 --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]

xlauko
xlauko approved these changes Jun 4, 2026
View reviewed changes
Copy link
Copy Markdown
Contributor

@xlauko xlauko left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

lgtm

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@xlauko xlauko xlauko approved these changes

@erichkeane erichkeane Awaiting requested review from erichkeane

@bcardosolopes bcardosolopes Awaiting requested review from bcardosolopes bcardosolopes is a code owner

@andykaylor andykaylor Awaiting requested review from andykaylor andykaylor is a code owner

Assignees

No one assigned

Labels

clang Clang issues not falling into any other category ClangIR Anything related to the ClangIR project

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

2 participants

@jsjodin @xlauko