//go:noescape
func gcmInit(productTable *[256]byte, h []byte)

//go:noescape
func gcmHash(output []byte, productTable *[256]byte, inp []byte, len int)

//go:noescape
func gcmMul(output []byte, productTable *[256]byte)

const (
	gcmCounterSize       = 16
	gcmBlockSize         = 16
	gcmTagSize           = 16
	gcmStandardNonceSize = 12
)

	VXOR XH3, XH3, XH3

	JMP tail_4x

done_4x:
#ifdef GOARCH_ppc64le
	VPERM   XL, XL, LEMASK, XL
#endif
	STXVD2X VXL, (XIP+R0)      // write out Xi
	RET

// func gcmMul(output []byte, productTable *[256]byte)
TEXT ·gcmMul(SB), NOSPLIT, $0-32
	MOVD output+0(FP), XIP
	MOVD productTable+24(FP), HTBL

	MOVD   $0x10, R8
	MOVD   $0x20, R9
	MOVD   $0x30, R10
	LXVD2X (XIP)(R0), VIN // load Xi

		uv := get2(U, V)
		if end(uv[0].(*rat)) != 0 || end(uv[1].(*rat)) != 0 {
			Z.dat <- finis
		} else {
			Z.dat <- mul(uv[0].(*rat), uv[1].(*rat))
			UU := Split(U)
			VV := Split(V)
			W := Add(Cmul(uv[0].(*rat), VV[0]), Cmul(uv[1].(*rat), UU[0]))
			<-Z.req
			Z.dat <- get(W)
			copy(Add(W, Mul(UU[1], VV[1])), Z)
		}
	}(U, V, Z)
	return Z
}

// Differentiate

func Diff(U PS) PS {
	Z := mkPS()

// Multiply a power series by a constant
func Cmul(c rat, U PS) PS {
	Z := mkPS()
	go func() {
		done := false
		for !done {
			<-Z.req
			u := get(U)
			if end(u) != 0 {
				done = true
			} else {
				Z.dat <- mul(c, u)
			}
		}
		Z.dat <- finis
	}()
	return Z
}

// Subtract

func Sub(U, V PS) PS {
	return Add(U, Cmul(neg(one), V))
}

    // CHECK: 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>
    // CHECK-NEXT: %[[convert_init:.*]] = mhlo.convert %[[init]] : tensor<f32>

  // 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>
  // CHECK-NEXT: %[[red1:.*]] = mhlo.reduce(%[[cmul]] init: %[[init]]) applies mhlo.add across dimensions = [0, 1, 2] : (tensor<8x8x8x8xf32>, tensor<f32>) -> tensor<8xf32>