{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)

  %0 = "tfl.dequantize"(%arg0) : (tensor<2x3x5x!quant.uniform<u8:f32, 0.1>>) -> tensor<2x3x5xf32>
  %1 = "tfl.slice"(%0, %arg1, %arg2) : (tensor<2x3x5xf32>, tensor<3xi32>, tensor<3xi32>) -> tensor<?x3x5xf32>
  func.return %1 : tensor<?x3x5xf32>

// CHECK: %0 = "tfl.dequantize"(%arg0)
// CHECK: %1 = "tfl.slice"(%0, %arg1, %arg2)
// CHECK: %2 = "tfl.quantize"(%1)

^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 private @"map/while_cond"(%arg0: tensor<i32>, %arg1: tensor<i32>, %arg2: tensor<!tf_type.variant<tensor<*xf32>>>, %arg3: tensor<?xf32>) -> tensor<i1> {
  %cst = "tf.Const"() {device = "/job:localhost/replica:0/task:0/device:CPU:0", value = dense<3> : tensor<i32>} : () -> tensor<i32>
  %0 = "tf.Less"(%arg0, %cst) {device = "/job:localhost/replica:0/task:0/device:CPU:0"} : (tensor<i32>, tensor<i32>) -> tensor<i1>

// CHECK-LABEL: func @reachable_function
func.func @reachable_function(%arg0: tensor<*x!tf_type.resource<tensor<2x8xi32>>>, %arg1: tensor<i32>) {
  // CHECK-NOT: "tf.AssignAddVariableOp"
  "tf.AssignAddVariableOp"(%arg0, %arg1) {dtype = "tfdtype$DT_INT32"} : (tensor<*x!tf_type.resource<tensor<2x8xi32>>>, tensor<i32>) -> ()
  "tf.OpA"(%arg0, %arg1) {f = @reachable_from_reachable_function} : (tensor<*x!tf_type.resource<tensor<2x8xi32>>>, tensor<i32>) -> ()
  func.return

// CHECK: %[[CONV:.*]] = stablehlo.convolution(%arg0, %arg1)
// CHECK: %[[MAX:.*]] = stablehlo.maximum %[[CONV]], %[[CONST]]
// CHECK: return %[[MAX]] : tensor<1x3x3x4xf32>
// CHECK: }

// -----

// CHECK-LABEL: @dot_general_with_relu_fn(
// CHECK-SAME:                 %[[ARG_0:.*]]: tensor<1x1x167xf32>,
func.func @dot_general_with_relu_fn(%arg0: tensor<1x1x167xf32>, %arg1: tensor<167x64xf32>) -> tensor<1x1x64xf32> {

// CHECK:       %0 = "tfl.while"(%arg0) ({
// CHECK:       ^bb0(%arg1: tensor<i32>):
// CHECK:         "tfl.yield"(%arg1) : (tensor<i32>) -> ()
// CHECK:       }, {
// CHECK:       ^bb0(%arg1: tensor<i32>):
// CHECK:         "tfl.yield"(%arg1) : (tensor<i32>) -> ()

// CHECK-LABEL: func @one_arg_per_TRI
func.func @one_arg_per_TRI(%arg0: tensor<i32>, %arg1: tensor<i32>) -> () {
  // CHECK: tf_device.replicate
  // CHECK-SAME: [%arg0, %arg1] as
  // CHECK-NOT: [%arg0, %arg1] as
  %ri = "tf.TPUReplicatedInput"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>

  func.func private @composite_add_fn(%arg0: tensor<1x2xf32>, %arg1: tensor<1x2xf32>) -> tensor<1x2xf32> attributes {_from_xla_call_module} {
    %0 = stablehlo.add %arg0, %arg1 : tensor<1x2xf32>
    return %0 : tensor<1x2xf32>
  }

		// ISEL  AuxInt ? arg0 : arg1
		// ISELZ is a special case of ISEL where arg1 is implicitly $0.
		//
		// AuxInt value indicates conditions 0=LT 1=GT 2=EQ 3=SO 4=GE 5=LE 6=NE 7=NSO.
		// ISEL accepts a CR bit argument, not a condition as expressed by AuxInt.
		// Convert the condition to a CR bit argument by the following conversion:
		//
		// AuxInt&3 ? arg0 : arg1 for conditions LT, GT, EQ, SO