ST (r)

	CALL p256SqrInternal(SB)	// rsqr = rˆ2
	ST (rsqr)

	LDacc (h)
	CALL p256SqrInternal(SB)	// hsqr = hˆ2
	ST (hsqr)

	LDt (h)
	CALL p256MulInternal(SB)	// hcub = hˆ3
	ST (hcub)

	LDt (y1in)
	CALL p256MulInternal(SB)	// y1 * hˆ3
	ST (s2)

	LDacc (x1in)
	LDt (hsqr)
	CALL p256MulInternal(SB)	// u1 * hˆ2
	ST (h)

	STx(r)

	CALL	p256SqrInternal<>(SB)    // rsqr = rˆ2
	STy	(rsqr)

	LDx(h)
	CALL	p256SqrInternal<>(SB)    // hsqr = hˆ2
	STy(hsqr)

	CALL	p256MulInternal<>(SB)    // hcub = hˆ3
	STy(hcub)

	LDx(y1in)
	CALL	p256MulInternal<>(SB)    // y1 * hˆ3
	STy(s2)

	LDP	hsqr(0*8), (x0, x1)
	LDP	hsqr(2*8), (x2, x3)
	LDP	0*16(a_ptr), (y0, y1)
	LDP	1*16(a_ptr), (y2, y3)

//              rsqrt(variance + epsilon)
// CHECK:  %[[RSQRT:.*]] = "tf.Rsqrt"(%[[ADD1]])
//              scale * rsqrt(variance + epsilon)
// CHECK:  %[[MUL1:.*]] = "tf.Mul"(%[[ARG1:.*]], %[[RSQRT]])
//              x * scale * rsqrt(variance + epsilon)
// CHECK:  %[[MUL2:.*]] = "tf.Mul"(%[[ARG0:.*]], %[[MUL1]])
//              mean * scale * rsqrt(variance + epsilon)

  // CHECK-DAG: %[[EPS_BCAST:.+]] = mhlo.constant dense<1.001000e-05> : tensor<256xf32>
  // CHECK-DAG: %[[VARIANCE_EPS:.+]] = mhlo.add %[[VARIANCE]], %[[EPS_BCAST]] : tensor<256xf32>
  // CHECK-DAG: %[[VARIANCE_EPS_RSQRT:.+]] = mhlo.rsqrt %[[VARIANCE_EPS]] : tensor<256xf32>
  // CHECK-DAG: %[[MULTIPLIER:.+]] = mhlo.multiply %[[VARIANCE_EPS_RSQRT]], %[[SCALE]] : tensor<256xf32>
  // CHECK-DAG: %[[MUL_MEAN:.+]] = mhlo.multiply %[[MULTIPLIER]], %[[MEAN]] : tensor<256xf32>

    $t, $m, $v, $beta, $gamma, F32Attr:$variance_epsilon, ConstBoolAttrFalse:$scale_after_normalization),
  (TF_AddOp
    (TF_MulOp $t, (TF_RsqrtOp:$rsqrt (TF_AddOp $v, (TF_ConstOp $variance_epsilon)))),
    (TF_SubOp $beta, (TF_MulOp $m, $rsqrt)))>;
def ConvertBatchNormWithGlobalNormalization_2 : Pat<
  (TF_BatchNormWithGlobalNormalizationOp

				// to check performance before making changes.
				zz = zz.sqr(z)
				zz, z = z, zz
				z = z.trunc(z, logM)

				zz = zz.sqr(z)
				zz, z = z, zz
				z = z.trunc(z, logM)

				zz = zz.sqr(z)
				zz, z = z, zz
				z = z.trunc(z, logM)

				zz = zz.sqr(z)
				zz, z = z, zz
				z = z.trunc(z, logM)
			}

	"⨒":                        "\u2a12",
	"⇉":                           "\u21c9",
	"›":                          "\u203a",
	"𝓇":                            "\U0001d4c7",
	"↱":                             "\u21b1",
	"]":                            "\u005d",
	"’":                           "\u2019",
	"’":                          "\u2019",

func (z *Rat) Mul(x, y *Rat) *Rat {
	if x == y {
		// a squared Rat is positive and can't be reduced (no need to call norm())
		z.a.neg = false
		z.a.abs = z.a.abs.sqr(x.a.abs)
		if len(x.b.abs) == 0 {
			z.b.abs = z.b.abs.setWord(1)
		} else {
			z.b.abs = z.b.abs.sqr(x.b.abs)
		}
		return z
	}
	z.a.Mul(&x.a, &y.a)
	z.b.abs = mulDenom(z.b.abs, x.b.abs, y.b.abs)
	return z.norm()
}

				}
			}
		}
	}
}

func testSqr(t *testing.T, x nat) {
	got := make(nat, 2*len(x))
	want := make(nat, 2*len(x))
	got = got.sqr(x)
	want = want.mul(x, x)
	if got.cmp(want) != 0 {
		t.Errorf("basicSqr(%v), got %v, want %v", x, got, want)
	}
}

func TestSqr(t *testing.T) {
	for _, a := range prodNN {
		if a.x != nil {

//CHECK-NEXT: return %0 : tensor<1x30x1xi32>
//CHECK-NEXT:}

func.func @rsqrt(%arg0: tensor<1x1x1x96xf32>) -> tensor<1x1x1x96xf32> {
 %0 = "vhlo.rsqrt_v1" (%arg0) : (tensor<1x1x1x96xf32>) -> tensor<1x1x1x96xf32>
 func.return %0 : tensor<1x1x1x96xf32>
}

//CHECK:func.func private @rsqrt(%arg0: tensor<1x1x1x96xf32>) -> tensor<1x1x1x96xf32> {