Integrate LLVM at llvm/llvm-project@799e9053641a (#2765)

Author: mrguenther

WORKSPACE.bazel

@@ -17,9 +17,9 @@ workspace(name = "stablehlo")
 
 load("@bazel_tools//tools/build_defs/repo:http.bzl", "http_archive")
 
-LLVM_COMMIT = "ac9049df7e62e2ca4dc5d103593b51639b5715e3"
+LLVM_COMMIT = "799e9053641a6478d3144866a97737b37b87c260"
 
-LLVM_SHA256 = "ea890ee3c13d9b2d70a359299a0b810c8bae9c729c5a94d81f5b304bf26f34b6"
+LLVM_SHA256 = "be33f1f9f20da6bd744d62356bf469e906e3b5f5e9cba2af6ee6418cee49f1f3"
 
 http_archive(
     name = "llvm-raw",

build_tools/llvm_version.txt

@@ -1 +1 @@
-ac9049df7e62e2ca4dc5d103593b51639b5715e3
+799e9053641a6478d3144866a97737b37b87c260

docs/generated/chlo.md

@@ -1968,14 +1968,6 @@ Syntax:
 >
 ```
 
-Enum cases:
-* EQ (`EQ`)
-* NE (`NE`)
-* GE (`GE`)
-* GT (`GT`)
-* LE (`LE`)
-* LT (`LT`)
-
 #### Parameters:
 
 | Parameter | C++ type | Description |
@@ -1994,13 +1986,6 @@ Syntax:
 >
 ```
 
-Enum cases:
-* NOTYPE (`NOTYPE`)
-* FLOAT (`FLOAT`)
-* TOTALORDER (`TOTALORDER`)
-* SIGNED (`SIGNED`)
-* UNSIGNED (`UNSIGNED`)
-
 #### Parameters:
 
 | Parameter | C++ type | Description |
@@ -2019,11 +2004,6 @@ Syntax:
 >
 ```
 
-Enum cases:
-* DEFAULT (`DEFAULT`)
-* HIGH (`HIGH`)
-* HIGHEST (`HIGHEST`)
-
 #### Parameters:
 
 | Parameter | C++ type | Description |

docs/generated/stablehlo_passes.md

@@ -259,6 +259,26 @@ func.func @add(%arg0: tensor<!quant.uniform<i8:f32, 1.000000e+00>>, %arg1: tenso
 
 _Legalize StableHLO to VHLO._
 
+Legalize StableHLO to the latest version of ops in VHLO. These ops can then
+be downgraded to older versions of VHLO for forward compatibility using
+`VhloToVersionPass`.
+
+```mlir
+stablehlo.exponential %[[ARG0]] <{result_accuracy = DEFAULT}> : tensor<f32>
+# ====>
+"vhlo.exponential_v2"(%[[ARG0]]) <{result_accuracy = #vhlo.DEFAULT_v1}> : !vhlo.tensor_v1<!vhlo.f32_v1>
+```
+
+See [vhlo.md > The VHLO dialect](https://github.com/openxla/stablehlo/blob/main/docs/vhlo.md)
+for full details on how VHLO is used to preserve forward and backward
+compatibility.
+
+#### Options
+
+```
+-allow-other-dialects : Allow serialization to use other (potentially unstable) dialects, inserts unrealized casts between dialects.
+```
+
 ### `-stablehlo-refine-arguments`
 
 _Refines the argument shapes of the main function._
@@ -279,6 +299,7 @@ func.func public @main(%arg0: tensor<16xf32>) -> tensor<?xf32> {
   %0 = stablehlo.custom_call @stablehlo.shape_refinement_operand_wrapper(%arg0, %c) {...}
     : (tensor<16xf32>, tensor<1xi64>) -> tensor<?xf32>
   ...
+}
 ```
 
 The `refinedTypesOption` can be used to specify a list of refined types.
@@ -459,7 +480,22 @@ _Legalize VHLO to StableHLO._
 
 ### `-vhlo-to-version`
 
-_Convert between versions of VHLO._
+_Convert between versions of VHLO for compatibility._
+
+Converts between versions of VHLO for IR upgrading and downgrading to
+preserve forward and backward compatibility.
+
+```mlir
+"vhlo.exponential_v2"(%[[ARG0]]) <{result_accuracy = DEFAULT}>
+# ==( -target=1.0.0 )==>
+"vhlo.exponential_v1"(%[[ARG0]])
+# ==( -target=1.9.0 )==>
+"vhlo.exponential_v2"(%[[ARG0]]) <{result_accuracy = DEFAULT}>
+```
+
+See [vhlo.md > The VHLO dialect](https://github.com/openxla/stablehlo/blob/main/docs/vhlo.md)
+for full details on how VHLO is used to preserve forward and backward
+compatibility.
 
 #### Options
 

stablehlo/dialect/Serialization.cpp

@@ -32,20 +32,25 @@ limitations under the License.
 #include "stablehlo/dialect/VhloOps.h"
 #include "stablehlo/transforms/Passes.h"
 
-#define DEBUG_TYPE "compat-passes"
+#define DEBUG_TYPE "stablehlo-compat"
 
 namespace mlir {
 namespace stablehlo {
 
 LogicalResult serializePortableArtifact(ModuleOp module,
                                         StringRef targetVersion,
-                                        raw_ostream& os) {
+                                        raw_ostream& os,
+                                        bool allowOtherDialects) {
   MLIRContext* context = module.getContext();
 
-  // Convert StableHLO --> VHLO. Will fail if entire program is not StableHLO.
+  // Convert StableHLO --> VHLO.
+  // If allowOtherDialects is true, we will allow other dialects to be present
+  // in the module, otherwise will fail if there are any other dialects present.
   {
     PassManager pm(context);
-    pm.addPass(stablehlo::createStablehloLegalizeToVhloPass());
+    StablehloLegalizeToVhloPassOptions options;
+    options.allowOtherDialects = allowOtherDialects;
+    pm.addPass(stablehlo::createStablehloLegalizeToVhloPass(options));
     if (!succeeded(pm.run(module))) {
       return failure();
     }

stablehlo/dialect/Serialization.h

@@ -34,7 +34,8 @@ namespace stablehlo {
 // unsupported dialects.
 LogicalResult serializePortableArtifact(ModuleOp module,
                                         StringRef targetVersion,
-                                        raw_ostream& os);
+                                        raw_ostream& os,
+                                        bool allowOtherDialects = false);
 
 // Read StableHLO portable artifact
 //

stablehlo/dialect/VhloTypes.cpp

@@ -322,6 +322,20 @@ void VhloTypeConverter::addVhloToBuiltinConversions() {
   });
 }
 
+namespace {
+Value materializeIllegalCast(OpBuilder& builder, Type type, ValueRange inputs,
+                             Location loc) {
+  return builder.create<UnrealizedConversionCastOp>(loc, type, inputs)
+      ->getResult(0);
+}
+}  // namespace
+
+void VhloTypeConverter::addUnrealizedMaterializations() {
+  addTargetMaterialization(materializeIllegalCast);
+  addSourceMaterialization(materializeIllegalCast);
+  addArgumentMaterialization(materializeIllegalCast);
+}
+
 namespace {
 // Helper functions for VHLO verifiers
 template <typename TypeOrAttr>

stablehlo/dialect/VhloTypes.h

@@ -55,6 +55,9 @@ class VhloTypeConverter : public VhloTypeConverterBase {
   // it is likely that a subclass should call this last, especially if a default
   // `Type -> Type` fallback conversion is registered.
   void addBuiltinToVhloConversions();
+
+  // Mark unrealized casts as legal. Useful for dialect mixing.
+  void addUnrealizedMaterializations();
 };
 
 // Autogenerated VHLO type printers and parsers.

stablehlo/tests/ops_stablehlo_quantized.mlir

@@ -1338,15 +1338,15 @@ func.func @quantized_element_type_c8(%arg0: tensor<1x2x!quant.uniform<i8<-128:12
 
 // -----
 
-// expected-error@+1 {{scale out of expressed type range}}
+// expected-error@+1 {{scale 1.055040e+05 out of expressed type range}}
 func.func @quantized_element_type_c6(%arg0: tensor<1x2x!quant.uniform<i4:f16, 10.550400e+04>>) {
    %0 = stablehlo.add %arg0,  %arg0 : tensor<1x2x!quant.uniform<i4:f16, 10.550400e+04>>
    func.return
 }
 
 // -----
 
-// expected-error@+1 {{scale out of expressed type range}}
+// expected-error@+1 {{scale 4.960464e-08 out of expressed type range}}
 func.func @quantized_element_type_c6(%arg0: tensor<1x2x!quant.uniform<i4:f16, 4.960464e-08>>) {
    %0 = stablehlo.add %arg0,  %arg0 : tensor<1x2x!quant.uniform<i4:f16, 4.960464e-08>>
    func.return

stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo_mixed.mlir

@@ -0,0 +1,189 @@
+// NOTE: Assertions have been autogenerated by utils/generate-test-checks.py
+
+// The script is designed to make adding checks to
+// a test case fast, it is *not* designed to be authoritative
+// about what constitutes a good test! The CHECK should be
+// minimized and named to reflect the test intent.
+
+// NOTE: Assertions have been autogenerated by utils/generate-test-checks.py
+
+// The script is designed to make adding checks to
+// a test case fast, it is *not* designed to be authoritative
+// minimized and named to reflect the test intent.
+
+// RUN: stablehlo-opt %s --stablehlo-legalize-to-vhlo=allow-other-dialects | FileCheck %s
+// RUN: stablehlo-opt %s > %t.0
+// RUN: stablehlo-opt %s --stablehlo-legalize-to-vhlo=allow-other-dialects | stablehlo-opt --vhlo-legalize-to-stablehlo > %t.1
+// RUN: diff %t.0 %t.1
+
+// CHECK-LABEL:   vhlo.func_v1 @op_other(
+// CHECK-SAME:                           %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<!vhlo.f32_v1>) -> (!vhlo.tensor_v1<!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_1:.*]] = builtin.unrealized_conversion_cast %[[VAL_0]] : !vhlo.tensor_v1<!vhlo.f32_v1> to tensor<f32>
+// CHECK:           %[[VAL_2:.*]] = arith.addf %[[VAL_1]], %[[VAL_1]] : tensor<f32>
+// CHECK:           %[[VAL_3:.*]] = builtin.unrealized_conversion_cast %[[VAL_2]] : tensor<f32> to !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_3]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @op_other(%arg0: tensor<f32>) -> tensor<f32> {
+  %0 = arith.addf %arg0, %arg0 : tensor<f32>
+  return %0 : tensor<f32>
+}
+
+// -----
+
+// CHECK-LABEL:   vhlo.func_v1 @op_shlo(
+// CHECK-SAME:                          %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<!vhlo.f32_v1>) -> (!vhlo.tensor_v1<!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_1:.*]] = "vhlo.add_v1"(%[[VAL_0]], %[[VAL_0]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>, !vhlo.tensor_v1<!vhlo.f32_v1>) -> !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_1]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @op_shlo(%arg0: tensor<f32>) -> tensor<f32> {
+  %0 = stablehlo.add %arg0, %arg0 : tensor<f32>
+  return %0 : tensor<f32>
+}
+
+// -----
+
+// CHECK-LABEL:   vhlo.func_v1 @mixed_shlo_other_shlo(
+// CHECK-SAME:                                        %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<!vhlo.f32_v1>) -> (!vhlo.tensor_v1<!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_1:.*]] = builtin.unrealized_conversion_cast %[[VAL_0]] : !vhlo.tensor_v1<!vhlo.f32_v1> to tensor<f32>
+// CHECK:           %[[VAL_2:.*]] = "vhlo.abs_v1"(%[[VAL_0]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>) -> !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           %[[VAL_3:.*]] = builtin.unrealized_conversion_cast %[[VAL_2]] : !vhlo.tensor_v1<!vhlo.f32_v1> to tensor<f32>
+// CHECK:           %[[VAL_4:.*]] = arith.addf %[[VAL_3]], %[[VAL_1]] : tensor<f32>
+// CHECK:           %[[VAL_5:.*]] = builtin.unrealized_conversion_cast %[[VAL_4]] : tensor<f32> to !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           %[[VAL_6:.*]] = "vhlo.abs_v1"(%[[VAL_5]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>) -> !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_6]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @mixed_shlo_other_shlo(%arg0: tensor<f32>) -> tensor<f32> {
+  %0 = stablehlo.abs %arg0 : tensor<f32>
+  %1 = arith.addf %0, %arg0 : tensor<f32>
+  %2 = stablehlo.abs %1 : tensor<f32>
+  return %2 : tensor<f32>
+}
+
+// -----
+
+// CHECK-LABEL:   vhlo.func_v1 @mixed_other_shlo_other(
+// CHECK-SAME:                                         %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<!vhlo.f32_v1>) -> (!vhlo.tensor_v1<!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_1:.*]] = builtin.unrealized_conversion_cast %[[VAL_0]] : !vhlo.tensor_v1<!vhlo.f32_v1> to tensor<f32>
+// CHECK:           %[[VAL_2:.*]] = arith.addf %[[VAL_1]], %[[VAL_1]] : tensor<f32>
+// CHECK:           %[[VAL_3:.*]] = builtin.unrealized_conversion_cast %[[VAL_2]] : tensor<f32> to !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           %[[VAL_4:.*]] = "vhlo.add_v1"(%[[VAL_3]], %[[VAL_0]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>, !vhlo.tensor_v1<!vhlo.f32_v1>) -> !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           %[[VAL_5:.*]] = builtin.unrealized_conversion_cast %[[VAL_4]] : !vhlo.tensor_v1<!vhlo.f32_v1> to tensor<f32>
+// CHECK:           %[[VAL_6:.*]] = arith.addf %[[VAL_5]], %[[VAL_1]] : tensor<f32>
+// CHECK:           %[[VAL_7:.*]] = builtin.unrealized_conversion_cast %[[VAL_6]] : tensor<f32> to !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_7]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @mixed_other_shlo_other(%arg0: tensor<f32>) -> tensor<f32> {
+  %0 = arith.addf %arg0, %arg0 : tensor<f32>
+  %1 = stablehlo.add %0, %arg0 : tensor<f32>
+  %2 = arith.addf %1, %arg0 : tensor<f32>
+  return %2 : tensor<f32>
+}
+
+// -----
+
+// CHECK-LABEL:   vhlo.func_v1 @op_with_region(
+// CHECK-SAME:                                 %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<1x16x16x320x!vhlo.f32_v1>,
+// CHECK-SAME:                                 %[[VAL_1:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<!vhlo.f32_v1>) -> (!vhlo.tensor_v1<1x320x!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_2:.*]] = "vhlo.reduce_v1"(%[[VAL_0]], %[[VAL_1]]) <{dimensions = #{{.*}}<dense<[1, 2]> : tensor<2xi64>>}> ({
+// CHECK:           ^bb0(%[[VAL_3:.*]]: !vhlo.tensor_v1<!vhlo.f32_v1>, %[[VAL_4:.*]]: !vhlo.tensor_v1<!vhlo.f32_v1>):
+// CHECK:             %[[VAL_5:.*]] = "vhlo.add_v1"(%[[VAL_3]], %[[VAL_4]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>, !vhlo.tensor_v1<!vhlo.f32_v1>) -> !vhlo.tensor_v1<!vhlo.f32_v1>
+// CHECK:             "vhlo.return_v1"(%[[VAL_5]]) : (!vhlo.tensor_v1<!vhlo.f32_v1>) -> ()
+// CHECK:           }) : (!vhlo.tensor_v1<1x16x16x320x!vhlo.f32_v1>, !vhlo.tensor_v1<!vhlo.f32_v1>) -> !vhlo.tensor_v1<1x320x!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_2]]) : (!vhlo.tensor_v1<1x320x!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @op_with_region(%arg0: tensor<1x16x16x320xf32>, %arg1: tensor<f32>) -> tensor<1x320xf32> {
+  %0 = stablehlo.reduce(%arg0 init: %arg1) applies stablehlo.add across dimensions = [1, 2] : (tensor<1x16x16x320xf32>, tensor<f32>) -> tensor<1x320xf32>
+  return %0 : tensor<1x320xf32>
+}
+
+// -----
+
+// CHECK-LABEL:   vhlo.func_v1 @op_with_region_mixed_other_shlo_other(
+// CHECK-SAME:                                                        %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<7x5x!vhlo.f32_v1>,
+// CHECK-SAME:                                                        %[[VAL_1:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<5x!vhlo.f32_v1>) -> (!vhlo.tensor_v1<5x!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_2:.*]] = "vhlo.reduce_v1"(%[[VAL_0]], %[[VAL_1]]) <{dimensions = #{{.*}}<dense<0> : tensor<1xi64>>}> ({
+// CHECK:           ^bb0(%[[VAL_3:.*]]: !vhlo.tensor_v1<5x!vhlo.f32_v1>, %[[VAL_4:.*]]: !vhlo.tensor_v1<5x!vhlo.f32_v1>):
+// CHECK:             %[[VAL_5:.*]] = builtin.unrealized_conversion_cast %[[VAL_4]] : !vhlo.tensor_v1<5x!vhlo.f32_v1> to tensor<5xf32>
+// CHECK:             %[[VAL_6:.*]] = builtin.unrealized_conversion_cast %[[VAL_3]] : !vhlo.tensor_v1<5x!vhlo.f32_v1> to tensor<5xf32>
+// CHECK:             %[[VAL_7:.*]] = arith.addf %[[VAL_6]], %[[VAL_5]] : tensor<5xf32>
+// CHECK:             %[[VAL_8:.*]] = builtin.unrealized_conversion_cast %[[VAL_7]] : tensor<5xf32> to !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:             %[[VAL_9:.*]] = "vhlo.add_v1"(%[[VAL_8]], %[[VAL_3]]) : (!vhlo.tensor_v1<5x!vhlo.f32_v1>, !vhlo.tensor_v1<5x!vhlo.f32_v1>) -> !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:             %[[VAL_10:.*]] = builtin.unrealized_conversion_cast %[[VAL_9]] : !vhlo.tensor_v1<5x!vhlo.f32_v1> to tensor<5xf32>
+// CHECK:             %[[VAL_11:.*]] = arith.addf %[[VAL_10]], %[[VAL_5]] : tensor<5xf32>
+// CHECK:             %[[VAL_12:.*]] = builtin.unrealized_conversion_cast %[[VAL_11]] : tensor<5xf32> to !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:             "vhlo.return_v1"(%[[VAL_12]]) : (!vhlo.tensor_v1<5x!vhlo.f32_v1>) -> ()
+// CHECK:           }) : (!vhlo.tensor_v1<7x5x!vhlo.f32_v1>, !vhlo.tensor_v1<5x!vhlo.f32_v1>) -> !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_2]]) : (!vhlo.tensor_v1<5x!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @op_with_region_mixed_other_shlo_other(%arg0: tensor<7x5xf32>, %arg1: tensor<5xf32>) -> tensor<5xf32> {
+  %0 = stablehlo.reduce(%arg0 init: %arg1) across dimensions = [0] : (tensor<7x5xf32>, tensor<5xf32>) -> tensor<5xf32>
+    reducer(%arg2: tensor<5xf32>, %arg3: tensor<5xf32>)  {
+    %1 = arith.addf %arg2, %arg3 : tensor<5xf32>
+    %2 = stablehlo.add %1, %arg2 : tensor<5xf32>
+    %3 = arith.addf %2, %arg3 : tensor<5xf32>
+    stablehlo.return %3 : tensor<5xf32>
+  }
+  return %0 : tensor<5xf32>
+}
+
+// -----
+
+// CHECK-LABEL:   vhlo.func_v1 @op_with_region_mixed_shlo_other_shlo(
+// CHECK-SAME:                                                       %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<7x5x!vhlo.f32_v1>,
+// CHECK-SAME:                                                       %[[VAL_1:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<5x!vhlo.f32_v1>) -> (!vhlo.tensor_v1<5x!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_2:.*]] = "vhlo.reduce_v1"(%[[VAL_0]], %[[VAL_1]]) <{dimensions = #{{.*}}<dense<0> : tensor<1xi64>>}> ({
+// CHECK:           ^bb0(%[[VAL_3:.*]]: !vhlo.tensor_v1<5x!vhlo.f32_v1>, %[[VAL_4:.*]]: !vhlo.tensor_v1<5x!vhlo.f32_v1>):
+// CHECK:             %[[VAL_5:.*]] = builtin.unrealized_conversion_cast %[[VAL_4]] : !vhlo.tensor_v1<5x!vhlo.f32_v1> to tensor<5xf32>
+// CHECK:             %[[VAL_6:.*]] = "vhlo.abs_v1"(%[[VAL_3]]) : (!vhlo.tensor_v1<5x!vhlo.f32_v1>) -> !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:             %[[VAL_7:.*]] = builtin.unrealized_conversion_cast %[[VAL_6]] : !vhlo.tensor_v1<5x!vhlo.f32_v1> to tensor<5xf32>
+// CHECK:             %[[VAL_8:.*]] = arith.addf %[[VAL_7]], %[[VAL_5]] : tensor<5xf32>
+// CHECK:             %[[VAL_9:.*]] = builtin.unrealized_conversion_cast %[[VAL_8]] : tensor<5xf32> to !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:             %[[VAL_10:.*]] = "vhlo.abs_v1"(%[[VAL_9]]) : (!vhlo.tensor_v1<5x!vhlo.f32_v1>) -> !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:             "vhlo.return_v1"(%[[VAL_10]]) : (!vhlo.tensor_v1<5x!vhlo.f32_v1>) -> ()
+// CHECK:           }) : (!vhlo.tensor_v1<7x5x!vhlo.f32_v1>, !vhlo.tensor_v1<5x!vhlo.f32_v1>) -> !vhlo.tensor_v1<5x!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_2]]) : (!vhlo.tensor_v1<5x!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @op_with_region_mixed_shlo_other_shlo(%arg0: tensor<7x5xf32>, %arg1: tensor<5xf32>) -> tensor<5xf32> {
+  %0 = stablehlo.reduce(%arg0 init: %arg1) across dimensions = [0] : (tensor<7x5xf32>, tensor<5xf32>) -> tensor<5xf32>
+    reducer(%arg2: tensor<5xf32>, %arg3: tensor<5xf32>)  {
+    %1 = stablehlo.abs %arg2 : tensor<5xf32>
+    %2 = arith.addf %1, %arg3 : tensor<5xf32>
+    %3 = stablehlo.abs %2 : tensor<5xf32>
+    stablehlo.return %3 : tensor<5xf32>
+  }
+  return %0 : tensor<5xf32>
+}
+
+// -----
+
+// CHECK-LABEL:   vhlo.func_v1 @stablehlo_in_other_op_region(
+// CHECK-SAME:                                               %[[VAL_0:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.tensor_v1<2x!vhlo.f32_v1>,
+// CHECK-SAME:                                               %[[VAL_1:[0-9]+|[a-zA-Z$._-][a-zA-Z0-9$._-]*]]: !vhlo.index_v1) -> (!vhlo.tensor_v1<2x!vhlo.f32_v1>) {
+// CHECK:           %[[VAL_2:.*]] = builtin.unrealized_conversion_cast %[[VAL_0]] : !vhlo.tensor_v1<2x!vhlo.f32_v1> to tensor<2xf32>
+// CHECK:           %[[VAL_3:.*]] = arith.constant 0 : index
+// CHECK:           %[[VAL_4:.*]] = arith.constant 1 : index
+// CHECK:           %[[VAL_5:.*]] = arith.constant 2 : index
+// CHECK:           %[[VAL_6:.*]] = arith.constant 0.000000e+00 : f32
+// CHECK:           %[[VAL_7:.*]] = scf.for %[[VAL_8:.*]] = %[[VAL_3]] to %[[VAL_5]] step %[[VAL_4]] iter_args(%[[VAL_9:.*]] = %[[VAL_2]]) -> (tensor<2xf32>) {
+// CHECK:             %[[VAL_10:.*]] = tensor.insert %[[VAL_6]] into %[[VAL_9]]{{\[}}%[[VAL_8]]] : tensor<2xf32>
+// CHECK:             %[[VAL_11:.*]] = builtin.unrealized_conversion_cast %[[VAL_10]] : tensor<2xf32> to !vhlo.tensor_v1<2x!vhlo.f32_v1>
+// CHECK:             %[[VAL_12:.*]] = "vhlo.add_v1"(%[[VAL_11]], %[[VAL_11]]) : (!vhlo.tensor_v1<2x!vhlo.f32_v1>, !vhlo.tensor_v1<2x!vhlo.f32_v1>) -> !vhlo.tensor_v1<2x!vhlo.f32_v1>
+// CHECK:             %[[VAL_13:.*]] = builtin.unrealized_conversion_cast %[[VAL_12]] : !vhlo.tensor_v1<2x!vhlo.f32_v1> to tensor<2xf32>
+// CHECK:             scf.yield %[[VAL_13]] : tensor<2xf32>
+// CHECK:           }
+// CHECK:           %[[VAL_14:.*]] = builtin.unrealized_conversion_cast %[[VAL_7]] : tensor<2xf32> to !vhlo.tensor_v1<2x!vhlo.f32_v1>
+// CHECK:           "vhlo.return_v1"(%[[VAL_14]]) : (!vhlo.tensor_v1<2x!vhlo.f32_v1>) -> ()
+// CHECK:         }
+func.func @stablehlo_in_other_op_region(%arg0: tensor<2xf32>, %arg1: index) -> tensor<2xf32> {
+  %c0 = arith.constant 0 : index
+  %c1 = arith.constant 1 : index
+  %c2 = arith.constant 2 : index
+  %cst = arith.constant 0.0 : f32
+
+  %for = scf.for %i = %c0 to %c2 step %c1 iter_args(%arg2 = %arg0) -> tensor<2xf32> {
+    %new_out = tensor.insert %cst into %arg2[%i] : tensor<2xf32>
+    %new_out_add = stablehlo.add %new_out, %new_out : tensor<2xf32>
+    scf.yield %new_out_add : tensor<2xf32>
+  }
+  return %for : tensor<2xf32>
+}