%add, %add_control = tf_executor.island wraps "tf.Add"(%arg2, %arg2) : (tensor<f32>, tensor<f32>) -> tensor<f32>
    %mul, %mul_control = tf_executor.island wraps "tf.Mul"(%arg2, %arg2) : (tensor<f32>, tensor<f32>) -> tensor<f32>

		{name: "FNMSUBS", argLength: 3, reg: fp31, asm: "FNMSUBS"}, // -arg0 + (arg1 * arg2)
		{name: "FNMSUBD", argLength: 3, reg: fp31, asm: "FNMSUBD"}, // -arg0 + (arg1 * arg2)
		{name: "MADD", argLength: 3, reg: gp31, asm: "MADD"},       // +arg0 + (arg1 * arg2)
		{name: "MADDW", argLength: 3, reg: gp31, asm: "MADDW"},     // +arg0 + (arg1 * arg2), 32-bit
		{name: "MSUB", argLength: 3, reg: gp31, asm: "MSUB"},       // +arg0 - (arg1 * arg2)

// MixedPrecision-NEXT: %[[v:.*]] = tfl.add %[[dq_1]], %arg2
// MixedPrecision-NEXT: return %[[v:.*]]
}

// CHECK-LABEL: quantize_float_placeholder_only
func.func @quantize_float_placeholder_only(%arg0: tensor<f32>, %arg1: tensor<2x3xi32>, %arg2: tensor<2x3xf32>) -> (tensor<f32>, tensor<2x3xi32>, tensor<2x3xf32>) {
  func.return %arg0, %arg1, %arg2: tensor<f32>, tensor<2x3xi32>, tensor<2x3xf32>

// CHECK-SAME (%arg0: !mlrt.Future, %arg1: !mlrt.Promise, %arg2: tensor<i32>, %arg3: tensor<i32>, %arg4: tensor<i32>, %arg5: tensor<!tf_type.resource<tensor<3x1xf32>>>, %arg6: tensor<?x3xf32>, %arg7: tensor<?x4xf32>)
// CHECK-NEXT: [[cst_0:%.*]] = "tf.Const"
// CHECK-NEXT: [[cst_1:%.*]] = "tf.Const"
// CHECK-NEXT: [[loop_counter:%.*]] = "tf.AddV2"(%arg2, [[cst_1]])
// CHECK-NEXT: [[weight:%.*]] = "tf.ReadVariableOp"(%arg5)

^bb0(%arg0: tensor<256x32x32x3xf32>, %arg1: tensor<3x3x3x16xf32>, %arg2: tensor<256x3x32x32xf32>) :
   // OK
   %0 = "tf.Conv2D"(%arg0, %arg1) {T = "tfdtype$DT_FLOAT", data_format = "NHWC", dilations = [1, 2, 3, 1], padding = "SAME", strides = [1, 4, 5, 1]} : (tensor<256x32x32x3xf32>, tensor<3x3x3x16xf32>) -> tensor<256x8x7x16xf32>
   // Unsupported data format

func.func @decompose_scalar_updates(%arg0: tensor<!tf_type.resource<tensor<4x3xi32>>>, %arg1: tensor<2xi32>, %arg2: tensor<i32>) {
  "tf_device.cluster"() ({
    // CHECK-NOT: ResourceScatterUpdate
    // CHECK-NOT: BroadcastTo
    // CHECK: TensorScatterUpdate
    "tf.ResourceScatterUpdate"(%arg0, %arg1, %arg2) : (tensor<!tf_type.resource<tensor<4x3xi32>>>, tensor<2xi32>, tensor<i32>) -> ()
    tf_device.return
  }) : () -> ()

// CHECK:     %0 = "tfl.shape"(%arg0) {tac.device = "DARWINN", tac.inference_type = "FLOAT"} : (tensor<?x?xi32>) -> tensor<2xi32>

  %cst = "tf.Const"() {value = dense<true> : tensor<i1>} : () -> tensor<i1>
  %1:2 = "tf.WhileRegion"(%cst, %arg1) ({
  ^bb0(%arg2: tensor<i1>, %arg3: tensor<i64>):
    "tf.Yield"(%arg2) : (tensor<i1>) -> ()
  }, {
  ^bb0(%arg2: tensor<i1>, %arg3: tensor<i64>):

// CHECK: }

// CHECK-LABEL: private @composite_dot_general_with_bias_same_shape_fn_1
// CHECK: %[[DOT_GENERAL:.*]] = stablehlo.dot_general %arg0, %arg1
// CHECK: %[[ADD:.*]] = stablehlo.add %[[DOT_GENERAL]], %arg2
// CHECK: return %[[ADD]] : tensor<1x3xf32>
// CHECK: }

// -----

// CHECK-LABEL: @conv_with_bias_fn(
// CHECK-SAME:                    %[[ARG_0:.*]]: tensor<1x3x3x4xf32>

func.func @replication_with_model_parallelism(%arg0: !rtype, %arg1: !rtype, %arg2: !rtype, %arg3: !rtype) -> (tensor<10x3xf32>, tensor<f32>) {
  %0 = "tf.opA"() {is_stateless = true} : () -> tensor<i32>
  %1 = "tf.opB"() {is_stateless = true} : () -> tensor<i32>
  %2 = "tf.TPUReplicatedInput"(%arg0, %arg2) : (!rtype, !rtype) -> !rtype
  %3 = "tf.TPUReplicatedInput"(%arg1, %arg3) : (!rtype, !rtype) -> !rtype