%conv = stablehlo.convolution(%quant_arg, %quant_cst) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = [[0, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x4x3x!quant.uniform<i8:f32, 0.0039207626791561354:-128>>, tensor<2x3x3x2x!quant.uniform<i8<-127:127>:f32, 0.015>>) -> tensor<1x3x4x2x!quant.uniform<i32:f32, 5.8949912267181218E-5>>...

			expectedScores: []int64{0, 99, 98, 97},
		},
		{
			scores:         []int64{1, 1, 1, 1},
			expectedScores: []int64{100, 100, 100, 100},
		},
		{
			scores:         []int64{1000, 1, 1, 1},
			expectedScores: []int64{100, 0, 0, 0},
		},
		{
			reverse:        true,
			scores:         []int64{0, 1, 1, 1},
			expectedScores: []int64{100, 0, 0, 0},
		},
		{
			scores:         []int64{0, 0, 0, 0},

    %0 = stablehlo.convolution(%arg0, %arg1) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = [[0, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x4x3xf32>, tensor<2x3x3x2xf32>) -> tensor<1x3x4x2xf32>
    return %0 : tensor<1x3x4x2xf32>
  }
}

// CHECK-LABEL: quantize_conv_per_tensor

    %0 = stablehlo.convolution(%arg0, %arg1) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = [[0, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x4x3xf32>, tensor<2x3x3x2xf32>) -> tensor<1x3x4x2xf32>
    return %0 : tensor<1x3x4x2xf32>
  }
}

// CHECK-LABEL: quantize_conv_fn

	// loops, so it is not safe to walk them without special handling.
	if !fs.isReparseTagNameSurrogate() {
		if fs.FileAttributes&syscall.FILE_ATTRIBUTE_DIRECTORY != 0 {
			m |= ModeDir | 0111
		}

		switch fs.filetype {
		case syscall.FILE_TYPE_PIPE:
			m |= ModeNamedPipe
		case syscall.FILE_TYPE_CHAR:
			m |= ModeDevice | ModeCharDevice
		}
	}

  add.arg0_data()[0] = 111;
  EXPECT_TRUE(add.Run());
  EXPECT_EQ(add.error_msg(), "");
  EXPECT_EQ(add.result0(), 153);
  EXPECT_EQ(add.result0_data()[0], 153);
  EXPECT_EQ(add.result0_data(), add.results()[0]);

  const AddWithCkptComp& add_const = add;
  EXPECT_EQ(add_const.error_msg(), "");
  EXPECT_EQ(add_const.arg0(), 111);
  EXPECT_EQ(add_const.arg0_data()[0], 111);

func mul(x, y float64) float64 { return x * y }
func div(x, y float64) float64 { return x / y }

func TestFloat64(t *testing.T) {
	base := []float64{
		0,
		math.Copysign(0, -1),
		-1,
		1,
		math.NaN(),
		math.Inf(+1),
		math.Inf(-1),
		0.1,
		1.5,
		1.9999999999999998,     // all 1s mantissa
		1.3333333333333333,     // 1.010101010101...
		1.1428571428571428,     // 1.001001001001...

	// if r3 is not zero (failed) then branch to finish
	BYTE $0xB9; BYTE $0x02; BYTE $0x00; BYTE $0x33 // lbl1     ltgr  3,3
	BYTE $0xA7; BYTE $0x74; BYTE $0x00; BYTE $0x08 // brc   b'0111',lbl2

	// stomic store shunt address in R5 into CEECAADMC
	BYTE $0xE3; BYTE $0x52; BYTE $0x00; BYTE $0x00; BYTE $0x00; BYTE $0x24 // stg   5,0(2)

00000040  a6 d3 b1 c1 5d 86 51 aa  63 67 6b 6e cb 3b 5e 59  |....].Q.cgkn.;^Y|
00000050  02 c2 57 fd 37 39 1b 73  9a 61 b0 78 de e8 cc f8  |..W.79.s.a.x....|
00000060  b3 01 11 e5 e9 31 85 4d  fe 60 d4 12 70 71 64 45  |.....1.M.`..pqdE|
00000070  e8 7d fb be 5b 82 c0 c4  e1 57 09 2c f2 d7 a3 79  |.}..[....W.,...y|
00000080  1c 40 08 e1 e6 cd e2 3e  e7 55 da 14 03 03 00 01  |.@.....>.U......|

    %0 = stablehlo.convolution(%arg0, %arg1) dim_numbers = [b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f], window = {pad = [[0, 1], [1, 1]]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x4x3xf32>, tensor<2x3x3x2xf32>) -> tensor<1x3x4x2xf32>
    return %0 : tensor<1x3x4x2xf32>
  }
  // CHECK: func private @composite_conv_fn