[(HasRankOf<1> $add_rhs_value),
   (HasEqualElementSize<[-1], [0]> $conv_out, $add_rhs)], [], (addBenefit -1)>;

// Convert conv+sub+mul pattern to conv+mul+add.
// (conv - sub) * mul -> conv * mul + (-sub) * mul
//
// This is needed to support Conv+BatchNorm pattern from Jax models converted
// using jax2tf w/o native serialization. Note that Jax2tf patterns always

// CHECK: %[[DEQUANTIZE:.*]] = "quantfork.dcast"(%[[QUANTIZE]])
// CHECK: %[[CONV:.*]] = "tf.Conv2D"(%arg0, %[[DEQUANTIZE]])
// CHECK: return %[[CONV]]
}

// CHECK-LABEL: perChannelFakeQuantWithConv2D
func.func @perChannelFakeQuantWithConv2D(tensor<256x32x32x3xf32>) -> (tensor<256x8x7x16xf32>) {
^bb0(%arg: tensor<256x32x32x3xf32>) :

// CHECK: %[[cast:.*]] = "tf.Cast"(%[[cst]]) <{Truncate = false}> : (tensor<2x3x3x2xbf16>) -> tensor<2x3x3x2xf32>
// CHECK: %[[conv:.*]] = "tf.Conv2D"(%[[ARG0]], %[[cast]])
// CHECK: %[[identity:.*]] = "tf.IdentityN"(%[[conv]]) {device = ""} : (tensor<1x3x2x2xf32>) -> tensor<1x3x2x2xf32>
// CHECK: return %[[identity]] : tensor<1x3x2x2xf32>

// -----

// Tests that `tf.BatchFunction` is inlined.

	s [2]string
}

func f(a T) { // ERROR "live at entry to f: a$"
	var e interface{} // ERROR "stack object e interface \{\}$"
	func() {          // ERROR "live at entry to f.func1: &e a$"
		e = a.s // ERROR "live at call to convT: &e$" "stack object a T$"
	}()
	// Before the fix, both a and e were live at the previous line.
	_ = e

// CHECK: %[[CONV:.+]] = stablehlo.convolution(%[[ARG1]], %[[ARG2]])
// CHECK-SAME: (tensor<1x3x4x3xf32>, tensor<2x3x3x2x!quant.uniform<i8<-127:127>:f32, 0.0023622048182750312>>) -> tensor<1x3x4x2xf32>
// CHECK: return %[[CONV]]

// -----

// Test that per-channel weight-only quantized dot_general op is produced when

    func.return %7 : tensor<1x112x112x32x!quant.uniform<u8:f32, 1.0>>

// CHECK: %[[conv:.*]] = "tfl.conv_2d"(%arg0, %arg1, %arg2)
// CHECK-SAME: -> tensor<1x112x112x32x!quant.uniform<u8:f32, 0.0078431372549019607:128>>
// CHECK: %[[cst:.*]] = "tfl.pseudo_qconst"()
// CHECK: %[[add:.*]] = tfl.add(%[[conv]], %[[cst]])
// CHECK-SAME: -> tensor<1x112x112x32x!quant.uniform<u8:f32, 1.000000e+00>>
// CHECK: return %[[add]]
}

        def results2 = new Object()

        def store1 = factory.createCachedStore("conf1")
        def store1b = factory.createCachedStore("conf1")
        def store2 = factory.createCachedStore("conf2")

        expect:
        store1.load(() -> results1) == results1
        store1.load(() -> { assert false }) == results1
        store1b.load(() -> { assert false }) == results1

func.func @not_lift_float_conv(%arg0: tensor<1x3x4x512xf32>) -> (tensor<*xf32>) {
  %cst = "tf.Const"() {value = dense<3.000000e+00> : tensor<2x3x512x512xf32>} : () -> tensor<2x3x512x512xf32>
  %0 = "tf.Conv2D"(%arg0, %cst) {
    data_format = "NHWC", device = "", dilations = [1, 1, 1, 1], explicit_paddings = [],
    padding = "SAME", strides = [1, 1, 2, 1], use_cudnn_on_gpu = true

// CHECK: %[[CONV:.+]] = stablehlo.convolution(%[[ARG1]], %[[ARG2]])
// CHECK-SAME: (tensor<1x3x4x3xf32>, tensor<2x3x3x2x!quant.uniform<i8:f32, 6.000000e-03:-128>>) -> tensor<1x3x4x2xf32>

  %6 = "tf.Cast"(%5) {Truncate = false} : (tensor<1x100x100x1xi8>) -> tensor<1x100x100x1xf32>

  %7 = "tf.Sub"(%6, %cst_4) : (tensor<1x100x100x1xf32>, tensor<f32>) -> tensor<1x100x100x1xf32>
  %8 = "tf.Conv2D"(%7, %cst_3) {dilations = [1, 1, 1, 1], padding = "VALID", strides = [1, 1, 1, 1]} : (tensor<1x100x100x1xf32>, tensor<3x3x1x1xf32>) -> tensor<1x98x98x1xf32>