// We cannot fuse this tfl.mul into the preceding conv op because %cst2 is not broadcast-compatible to %cst0.
  %1 = "tfl.mul"(%0, %cst2) {fused_activation_function = "RELU6"} : (tensor<1x4x4x2xf32>, tensor<4x2xf32>) -> tensor<1x4x4x2xf32>

  func.return %1 : tensor<1x4x4x2xf32>

// CHECK:  %0 = "tfl.depthwise_conv_2d"(%arg0, %cst, %cst_0)
// CHECK:  %1 = tfl.mul(%0, %cst_1)
// CHECK:  return %1
}

	MUL:               "MUL",
	MUL_ZZ:            "MUL.ZZ",
	MUL_S_EQ:          "MUL.S.EQ",
	MUL_S_NE:          "MUL.S.NE",
	MUL_S_CS:          "MUL.S.CS",
	MUL_S_CC:          "MUL.S.CC",
	MUL_S_MI:          "MUL.S.MI",
	MUL_S_PL:          "MUL.S.PL",
	MUL_S_VS:          "MUL.S.VS",
	MUL_S_VC:          "MUL.S.VC",
	MUL_S_HI:          "MUL.S.HI",
	MUL_S_LS:          "MUL.S.LS",

  %178 = "tf.Pack"(%7, %177) {axis = 0 : i64, device = ""} : (tensor<i32>, tensor<i32>) -> tensor<2xi32>
  %179 = "tf.Tile"(%115, %178) {device = ""} : (tensor<?x1xi64>, tensor<2xi32>) -> tensor<?x?xi64>
  %180 = "tf.Mul"(%177, %118) {device = ""} : (tensor<i32>, tensor<i32>) -> tensor<i32>
  %181 = "tf.Pack"(%180) {axis = 0 : i64, device = ""} : (tensor<i32>) -> tensor<1xi32>

  // CHECK-NEXT: %[[sub:.*]] = mhlo.subtract %[[act]], %[[bcast_arg3]] : tensor<8x8x8x8xf32>
  // CHECK-NEXT: %[[mul:.*]] = mhlo.multiply %[[grad]], %[[sub]] : tensor<8x8x8x8xf32>
  // CHECK-NEXT: mhlo.constant dense<[0, 1, 2]> : tensor<3xi64>
  // CHECK-NEXT: %[[cmul:.*]] = mhlo.convert %[[mul]] : tensor<8x8x8x8xf32>
  // CHECK-NEXT: %[[init:.*]] = mhlo.constant dense<-0.000000e+00> : tensor<f32>

		}
		break
	}
	// match: (MUL _ (MOVWconst [0]))
	// result: (MOVWconst [0])
	for {
		for _i0 := 0; _i0 <= 1; _i0, v_0, v_1 = _i0+1, v_1, v_0 {
			if v_1.Op != OpARMMOVWconst || auxIntToInt32(v_1.AuxInt) != 0 {
				continue
			}
			v.reset(OpARMMOVWconst)
			v.AuxInt = int32ToAuxInt(0)
			return true
		}
		break
	}
	// match: (MUL x (MOVWconst [1]))
	// result: x
	for {

}

// CHECK-LABEL:   func @mul(
// CHECK-SAME:              %[[VAL_0:.*]]: tensor<2xi32>) -> tensor<2xi32> {
// CHECK:           %[[VAL_1:.*]] = "tf.Mul"(%[[VAL_0]], %[[VAL_0]]) : (tensor<2xi32>, tensor<2xi32>) -> tensor<2xi32>
// CHECK:           return %[[VAL_1]] : tensor<2xi32>
// CHECK:         }
func.func @mul(%arg0: tensor<2xi32>) -> tensor<2xi32> {

	/*13140*/ uint16(xCondDataSize), 13144, 13148, 0,
	/*13144*/ uint16(xSetOp), uint16(MUL),
	/*13146*/ uint16(xArgRM16),
	/*13147*/ uint16(xMatch),
	/*13148*/ uint16(xSetOp), uint16(MUL),
	/*13150*/ uint16(xArgRM32),
	/*13151*/ uint16(xMatch),
	/*13152*/ uint16(xCondDataSize), 13144, 13148, 13156,
	/*13156*/ uint16(xSetOp), uint16(MUL),
	/*13158*/ uint16(xArgRM64),
	/*13159*/ uint16(xMatch),

			loopHead := s.f.NewBlock(ssa.BlockPlain)
			loopBody := s.f.NewBlock(ssa.BlockPlain)

			// Pick right size ops.
			var mul, and, add, zext ssa.Op
			if s.config.PtrSize == 4 {
				mul = ssa.OpMul32
				and = ssa.OpAnd32
				add = ssa.OpAdd32
				zext = ssa.OpCopy
			} else {
				mul = ssa.OpMul64
				and = ssa.OpAnd64
				add = ssa.OpAdd64
				zext = ssa.OpZeroExt32to64
			}

    StringRef reduce_op = op.getReduceOp();
    if (reduce_op == "Add") {
      BuildReduceBody<AddOp>(element_type, &reduce_scatter.getComputation(),
                             &rewriter);
    } else if (reduce_op == "Mul") {
      BuildReduceBody<MulOp>(element_type, &reduce_scatter.getComputation(),
                             &rewriter);
    } else if (reduce_op == "Min") {

8�_�s�_:.�_��b����㘿??�st:�:��|.�����T�����n�ceK�����RE��UUUS�[�������UUU^n��r�h�`yY�i��h�`F��3����o��&�#\~��&[�}I����II�������"I�%�l߉4��H 4��Ka�L�a�� �F�&Ɖ�"bL��2��m63��[m�US32 Ń�̓l�2�����f�-��Hk�VX�K�SJRllfR�l�(�����Хk*��mUl�lE�i6K1-���V��d��+ZI��2��b�0,2+��*c b�C1 ��ш4�X��dJ���uT2K%�Y���j�6)���E���6�fX�ɪl,�,�iX�T6Z�͔�I���I� ؓ%���+k`Ya�3"$j�TT�I1�jj�3 f �,���Դ�V�i2��LE,McVa�m�R�1e�5�D�X� �ʬ���D�2̓&�4[$�� m�F���%�B��K��«L�4FZ6�2��m�ڴ���U4�FF��V�#,2�T��-�ٲ�Z$ҍ1...