{"Name":"CMHS (register)","Bits":"0|1|1|1|1|1|1|0|size:2|1|Rm:5|0|0|1|1|1|1|Rn:5|Rd:5","Arch":"Scalar Scalar variant","Syntax":"CMHS <V><d>, <V><n>, <V><m>","Code":"","Alias":""},
{"Name":"CMHS (register)","Bits":"0|Q|1|0|1|1|1|0|size:2|1|Rm:5|0|0|1|1|1|1|Rn:5|Rd:5","Arch":"Vector Vector variant","Syntax":"CMHS <Vd>.<T>, <Vn>.<T>, <Vm>.<T>","Code":"","Alias":""},

}

// from exp/4s/data.go
var pieces4 = []Piece{
	{0, 0, Point{4, 1}, []Point{{0, 0}, {1, 0}, {1, 0}, {1, 0}}, nil, nil},
	{1, 0, Point{1, 4}, []Point{{0, 0}, {0, 1}, {0, 1}, {0, 1}}, nil, nil},
	{2, 0, Point{4, 1}, []Point{{0, 0}, {1, 0}, {1, 0}, {1, 0}}, nil, nil},
	{3, 0, Point{1, 4}, []Point{{0, 0}, {0, 1}, {0, 1}, {0, 1}}, nil, nil},
}

var _ = [42]*T{
	{},
	{1, 2},
	{3, 4},
}

}

func ExampleTrailingZeros8() {
	fmt.Printf("TrailingZeros8(%08b) = %d\n", 14, bits.TrailingZeros8(14))
	// Output:
	// TrailingZeros8(00001110) = 1
}

func ExampleTrailingZeros16() {
	fmt.Printf("TrailingZeros16(%016b) = %d\n", 14, bits.TrailingZeros16(14))
	// Output:
	// TrailingZeros16(0000000000001110) = 1
}

func ExampleTrailingZeros32() {

)

const (
	t1 = 0x00 // 0000 0000
	tx = 0x80 // 1000 0000
	t2 = 0xC0 // 1100 0000
	t3 = 0xE0 // 1110 0000
	t4 = 0xF0 // 1111 0000
	t5 = 0xF8 // 1111 1000

	maskx = 0x3F // 0011 1111
	mask2 = 0x1F // 0001 1111
	mask3 = 0x0F // 0000 1111
	mask4 = 0x07 // 0000 0111

	rune1Max = 1<<7 - 1
	rune2Max = 1<<11 - 1
	rune3Max = 1<<16 - 1

	// The default lowest and highest continuation byte.
	locb = 0x80 // 1000 0000

	// and "1111" for SHAKE, or 00000010b and 00001111b, respectively. Then the
	// padding rule from section 5.1 is applied to pad the message to a multiple
	// of the rate, which involves adding a "1" bit, zero or more "0" bits, and
	// a final "1" bit. We merge the first "1" bit from the padding into dsbyte,
	// giving 00000110b (0x06) and 00011111b (0x1f).

    // CHECK: %[[CALL:.*]] = call @quantized_dot_general_fn_1(%[[Q2]], %[[Q1]])

    // CHECK: %[[REDUCE:.*]] = "stablehlo.reduce_window"(%[[CALL]], %[[Q0]])
    // CHECK{LITERAL}: padding = dense<[[0, 0], [1, 1], [1, 1], [0, 0]]> : tensor<4x2xi64>
    // CHECK-SAME: window_dimensions = array<i64: 1, 3, 3, 1>
    // CHECK: %[[ARG1:.*]]: tensor<!quant.uniform<i8:f32, 3.000000e-01:1>>, %[[ARG2:.*]]: tensor<!quant.uniform<i8:f32, 3.000000e-01:1>>

000000e0  00 00 00 00 00 94 68 2d  a3 82 51 ed 14 ef 68 ca  |......h-..Q...h.|
000000f0  42 c5 5c 95 7a 77 9a 7a  61 99 c6 44 1e e6 5d 71  |B.\.zw.za..D..]q|
00000100  41 c8 77 d1 d4 b6 39 bb  38 57 41 b8 41 f2 02 96  |A.w...9.8WA.A...|
00000110  d0 b0 5b 99 76 3c 2b 52  f2 8f 5f 33 94 fa 66 74  |..[.v<+R.._3..ft|
00000120  cc f0 60 1b ee 0a 38 30  78 b2 9e 54 16 d6 71 0e  |..`...80x..T..q.|
00000130  b2 ea 4e 0f 13 f0 6e 63  88 e7 9f 55 65 0b 00 00  |..N...nc...Ue...|

  %1 = "stablehlo.reduce_window"(%arg0, %0) ({
    ^bb0(%arg1: tensor<f32>, %arg2: tensor<f32>):
      %2 = stablehlo.maximum %arg1, %arg2 : tensor<f32>
      stablehlo.return %2 : tensor<f32>
  }) {padding = dense<[[0, 0], [1, 1], [1, 1], [0, 0]]> : tensor<4x2xi64>, window_dimensions = array<i64: 1, 3, 3, 1>} : (tensor<2x3x1x3xf32>, tensor<f32>) -> tensor<2x3x1x3xf32>
  return %1 : tensor<2x3x1x3xf32>
}

// -----

000000e0  37 ef 0d 19 51 88 35 75  71 b5 e5 54 5b 12 2e 8f  |7...Q.5uq..T[...|
000000f0  09 67 fd a7 24 20 3e b2  56 1c ce 97 28 5e f8 2b  |.g..$ >.V...(^.+|
00000100  2d 4f 9e f1 07 9f 6c 4b  5b 83 56 e2 32 42 e9 58  |-O....lK[.V.2B.X|
00000110  b6 d7 49 a6 b5 68 1a 41  03 56 6b dc 5a 89        |..I..h.A.Vk.Z.|
>>> Flow 2 (server to client)
00000000  16 03 03 00 9b 02 00 00  97 03 03 1c 1a ee 30 2f  |..............0/|

000000e0  45 00 17 00 41 04 1e 18  37 ef 0d 19 51 88 35 75  |E...A...7...Q.5u|
000000f0  71 b5 e5 54 5b 12 2e 8f  09 67 fd a7 24 20 3e b2  |q..T[....g..$ >.|
00000100  56 1c ce 97 28 5e f8 2b  2d 4f 9e f1 07 9f 6c 4b  |V...(^.+-O....lK|
00000110  5b 83 56 e2 32 42 e9 58  b6 d7 49 a6 b5 68 1a 41  |[.V.2B.X..I..h.A|
00000120  03 56 6b dc 5a 89                                 |.Vk.Z.|
>>> Flow 4 (server to client)