op_VME    uint32 = 0xE7A6 // 	VRR-c	VECTOR MULTIPLY EVEN
	op_VMH    uint32 = 0xE7A3 // 	VRR-c	VECTOR MULTIPLY HIGH
	op_VMLE   uint32 = 0xE7A4 // 	VRR-c	VECTOR MULTIPLY EVEN LOGICAL
	op_VMLH   uint32 = 0xE7A1 // 	VRR-c	VECTOR MULTIPLY HIGH LOGICAL
	op_VMLO   uint32 = 0xE7A5 // 	VRR-c	VECTOR MULTIPLY ODD LOGICAL
	op_VML    uint32 = 0xE7A2 // 	VRR-c	VECTOR MULTIPLY LOW
	op_VMO    uint32 = 0xE7A7 // 	VRR-c	VECTOR MULTIPLY ODD

(Div64 n (Const64 [-1<<63])) && isNonNegative(n)                 => (Const64 [0])

// Unsigned divide, not a power of 2.  Strength reduce to a multiply.
// For 8-bit divides, we just do a direct 9-bit by 8-bit multiply.
(Div8u x (Const8 [c])) && umagicOK8(c) =>
  (Trunc32to8
    (Rsh32Ux64 <typ.UInt32>
      (Mul32 <typ.UInt32>
        (Const32 <typ.UInt32> [int32(1<<8+umagic8(c).m)])

		// Let x = arg0*arg1 (full 32x32->64  unsigned multiply). Returns uint32(x), and flags set to overflow if uint32(x) != x.
		{name: "MULLU", argLength: 2, reg: regInfo{inputs: []regMask{ax, gpsp}, outputs: []regMask{ax, 0}, clobbers: dx}, typ: "(UInt32,Flags)", asm: "MULL", commutative: true, clobberFlags: true},
		// Let x = arg0*arg1 (full 64x64->128 unsigned multiply). Returns uint64(x), and flags set to overflow if uint64(x) != x.

	/* Vector multiply */
	{as: AVMULESB, a1: C_VREG, a2: C_VREG, a6: C_VREG, type_: 82, size: 4},              /* vector multiply, vx-form */
	{as: AVPMSUM, a1: C_VREG, a2: C_VREG, a6: C_VREG, type_: 82, size: 4},               /* vector polynomial multiply & sum, vx-form */
	{as: AVMSUMUDM, a1: C_VREG, a2: C_VREG, a3: C_VREG, a6: C_VREG, type_: 83, size: 4}, /* vector multiply-sum, va-form */

	/* Vector rotate */

int64_t L2NormalizationOp::GetArithmeticCount(Operation* op) {
  int64_t count;
  // Computing the squared L2 norm is N multiply-adds so 2N ops,
  // then the single inverse-sqrt is negligible, then we multiply each
  // value by the resulting multiplier, so an extra N ops. count 3N ops.
  if (ArithmeticCountUtilHelper::GetFirstOutputCount(op, &count)) {
    return 3 * count;
  }

  LogicalResult matchAndRewrite(TFL::MulOp mul_op,
                                PatternRewriter &rewriter) const override {
    // If we are broadcasting on the lhs then don't fold the multiply as it
    // would increase the amount of compute done by the fully connected op.
    if (mul_op.getLhs().getType() != mul_op.getType()) return failure();

    // Mul.
    DenseElementsAttr cst;

      mlir::cast<ShapedType>(dot_op.getResult().getType()), dot_op.getLoc());
}

// Converts mhlo.dot to tf.BatchMatMul. Reshape or Transpose ops will also be
// inserted to convert to well-formed matrix multiply.
Value ConvertDotGeneralOp(PatternRewriter& rewriter, Operation* old_op) {
  auto dot_general_op = cast<mhlo::DotGeneralOp>(old_op);
  return ConvertDot(
      rewriter, dot_general_op.getLhs(), dot_general_op.getRhs(),

     And<[TFL_OperandHasRankAtMostPred<0, 5>,
          TFL_OperandHasRankAtMostPred<1, 5>]>>,
   DynamicRangeQuantizedOpInterface]> {

  let summary = "Batch Matrix Multiply Operator";

  let description = [{
Performs a batched matrix multiplication on the inputs. Follows the
conventions of TensorFlow BatchMatMulV2, with support for unknown dimensions

	case _Gwaiting:
		if !gp.waitreason.isMutexWait() {
			// Not blocking on a lock.
			break
		}
		// Blocking on a lock, measure it. Note that because we're
		// sampling, we have to multiply by our sampling period to get
		// a more representative estimate of the absolute value.
		// gTrackingPeriod also represents an accurate sampling period
		// because we can only enter this state from _Grunning.

the precision of the target type, preventing fusion that would discard that rounding.
</p>

<p>
For instance, some architectures provide a "fused multiply and add" (FMA) instruction
that computes <code>x*y + z</code> without rounding the intermediate result <code>x*y</code>.
These examples show when a Go implementation can use that instruction:
</p>

<pre>