SBFX_ZZ:           "SBFX.ZZ",
	SDIV_EQ:           "SDIV.EQ",
	SDIV_NE:           "SDIV.NE",
	SDIV_CS:           "SDIV.CS",
	SDIV_CC:           "SDIV.CC",
	SDIV_MI:           "SDIV.MI",
	SDIV_PL:           "SDIV.PL",
	SDIV_VS:           "SDIV.VS",
	SDIV_VC:           "SDIV.VC",
	SDIV_HI:           "SDIV.HI",
	SDIV_LS:           "SDIV.LS",
	SDIV_GE:           "SDIV.GE",
	SDIV_LT:           "SDIV.LT",

	/*13209*/ uint16(xCondDataSize), 13213, 13217, 0,
	/*13213*/ uint16(xSetOp), uint16(IDIV),
	/*13215*/ uint16(xArgRM16),
	/*13216*/ uint16(xMatch),
	/*13217*/ uint16(xSetOp), uint16(IDIV),
	/*13219*/ uint16(xArgRM32),
	/*13220*/ uint16(xMatch),
	/*13221*/ uint16(xCondDataSize), 13213, 13217, 13225,
	/*13225*/ uint16(xSetOp), uint16(IDIV),
	/*13227*/ uint16(xArgRM64),
	/*13228*/ uint16(xMatch),

	DCTQPQ:         "dctqpq",
	DCTQPQCC:       "dctqpq.",
	DDEDPD:         "ddedpd",
	DDEDPDCC:       "ddedpd.",
	DDEDPDQ:        "ddedpdq",
	DDEDPDQCC:      "ddedpdq.",
	DDIV:           "ddiv",
	DDIVCC:         "ddiv.",
	DDIVQ:          "ddivq",
	DDIVQCC:        "ddivq.",
	DENBCD:         "denbcd",
	DENBCDCC:       "denbcd.",
	DENBCDQ:        "denbcdq",
	DENBCDQCC:      "denbcdq.",

// Fusing-LABEL: FusingdivRelu
// Fusing:  %[[div:[0-9].*]] = tfl.div %arg0, %arg1 {fused_activation_function = "NONE"} : tensor<1xf32>
// Fusing:  %[[div1:[0-9].*]] = tfl.div %arg0, %[[div]] {fused_activation_function = "RELU"} : tensor<1xf32>
// Fusing:  %[[relu:[0-9].*]] = "tfl.relu"(%arg0) : (tensor<1xf32>) -> tensor<1xf32>
// Fusing:  %[[div2:[0-9].*]] = tfl.div %[[relu]], %[[div1]] {fused_activation_function = "RELU6"} : tensor<1xf32>

// CHECK-LABEL: func @floordiv_broadcast_i32
func.func @floordiv_broadcast_i32(%arg0: tensor<2x3xi32>, %arg1: tensor<3xi32>) -> tensor<2x3xi32> {
  // CHECK-DAG: [[DIV:%.+]] = chlo.broadcast_divide %arg0, %arg1 {broadcast_dimensions = array<i64: 1>}
  // CHECK-DAG: [[MUL:%.+]] = chlo.broadcast_multiply [[DIV]], %arg1 {broadcast_dimensions = array<i64: 1>}

  func.return %0 : tensor<4x4x4x4xi32>
}

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

	{ir.ODIV, types.TFLOAT32}: ssa.OpDiv32F,
	{ir.ODIV, types.TFLOAT64}: ssa.OpDiv64F,

	{ir.ODIV, types.TINT8}:   ssa.OpDiv8,
	{ir.ODIV, types.TUINT8}:  ssa.OpDiv8u,
	{ir.ODIV, types.TINT16}:  ssa.OpDiv16,
	{ir.ODIV, types.TUINT16}: ssa.OpDiv16u,
	{ir.ODIV, types.TINT32}:  ssa.OpDiv32,
	{ir.ODIV, types.TUINT32}: ssa.OpDiv32u,
	{ir.ODIV, types.TINT64}:  ssa.OpDiv64,

ngth);if(n||a){e.remove();return}e.value=r}parse(e){let t=zm(e),r=[],n=[];for(let a of t.nodes)n.push(a),a.type==="div"&&a.value===","&&(r.push(n),n=[]);return r.push(n),r.filter(a=>a.length>0)}stringify(e){if(e.length===0)return"";let t=[];for(let r of e)r[r.length-1].type!=="div"&&r.push(this.div(e)),t=t.concat(r);return t[0].type==="div"&&(t=t.slice(1)),t[t.length-1].type==="div"&&(t=t.slice(0,-2+1||void 0)),zm.stringify({nodes:t})}clone(e,t,r){let n=[],a=!1;for(let s of r)!a&&s.type==="word"...

// division can result in strange behavior.
//
//      floordiv = cast(floordiv(cast(left), cast(right))))
//
//   %left_cast = cast(%left)
//   %right_cast = cast(%right)
//   %div = div(%left, %left)
//   %floored = floor(%div)
//   %floored_cast = cast(%floored)
//
// Required to manually specify the intermediate types.
class ConvertBF16FloorDivOp : public OpRewritePattern<TF::FloorDivOp> {
 public: