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)

  // 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>

			// V(k'+1) = V(2k+2) = V(k+1)² - 2.
			t1 = t1.sqr(vk1)
			t1 = t1.add(t1, nm2)
			t2, vk1 = t2.div(vk1, t1, n)
		} else {
			// k' = 2k
			// V(k'+1) = V(2k+1) = V(k) V(k+1) - P.
			t1 = t1.mul(vk, vk1)
			t1 = t1.add(t1, n)
			t1 = t1.sub(t1, natP)
			t2, vk1 = t2.div(vk1, t1, n)
			// V(k') = V(2k) = V(k)² - 2
			t1 = t1.sqr(vk)
			t1 = t1.add(t1, nm2)
			t2, vk = t2.div(vk, t1, n)
		}

    $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

}

TEST_F(CWiseUnaryGradTest, Rsqrt) {
  auto x_fn = [this](const int i) { return RV({1, 2, 3, 4, 5, 6, 7, 8}); };
  TestCWiseGrad<float, float>(RSQRT, x_fn);
}

TEST_F(CWiseUnaryGradTest, Rsqrt_Complex) {
  auto x_fn = [this](const int i) {
    return CRV({{-1.0f, 0.5f}, {1.0f, 0.5f}, {2, -1}});
  };
  TestCWiseGrad<complex64, complex64>(RSQRT, x_fn);
}

TEST_F(CWiseUnaryGradTest, Exp) {

				// 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)
			}

	x2sq.Square(&x)

	if x2 != x2sq {
		t.Fatalf("all ones failed\nmul: %x\nsqr: %x\n", x2, x2sq)
	}

	var bytes [32]byte

	_, err := io.ReadFull(rand.Reader, bytes[:])
	if err != nil {
		t.Fatal(err)
	}
	x.SetBytes(bytes[:])

	x2.Multiply(&x, &x)
	x2sq.Square(&x)

	if x2 != x2sq {
		t.Fatalf("all ones failed\nmul: %x\nsqr: %x\n", x2, x2sq)
	}
}

func TestEqual(t *testing.T) {

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 {