recurrent activation function for GRU from `hard_sigmoid` to `sigmoid`, and
    `reset_after` to True in 2.0. Historically recurrent activation is
    `hard_sigmoid` since it is fast than 'sigmoid'. With new unified backend
    between CPU and GPU mode, since the CuDNN kernel is using sigmoid, we change
    the default for CPU mode to sigmoid as well. With that, the default GRU will

  %0 = "tf.Neg"(%arg0) : (tensor<?xf32>) -> tensor<?xf32>
  func.return %0 : tensor<?xf32>
}

// -----

// CHECK-LABEL: @sigmoid
func.func @sigmoid(%arg0: tensor<2xf32>) -> tensor<2xf32> {
  // CHECK: mhlo.logistic
  %0 = "tf.Sigmoid"(%arg0) : (tensor<2xf32>) -> tensor<2xf32>
  func.return %0 : tensor<2xf32>
}

// -----

// CHECK-LABEL: @sigmoid_complex

            "Floor", "IsFinite", "IsInf", "IsNan", "Inv", "Reciprocal", "Log",
            "Log1p", "Invert", "LogicalNot", "Ndtri", "Neg", "Rint", "Round",
            "Rsqrt", "Sigmoid", "Sign", "Sinh", "Softplus", "Softsign", "Sqrt",
            "Square", "Tan", "Tanh", "Real", "Imag", "Erf", "Erfc", "Erfinv",
            "Lgamma", "Digamma",
            // Binary

	}

	var (
		err  error
		m, c *bigmod.Nat
		N    *bigmod.Modulus
		t0   = bigmod.NewNat()
	)
	if priv.Precomputed.n == nil {
		N, err = bigmod.NewModulusFromBig(priv.N)
		if err != nil {
			return nil, ErrDecryption
		}
		c, err = bigmod.NewNat().SetBytes(ciphertext, N)
		if err != nil {
			return nil, ErrDecryption
		}
		m = bigmod.NewNat().Exp(c, priv.D.Bytes(), N)
	} else {

// CHECK:         }
func.func @neg_dynamic(%arg0: tensor<?xf32>) -> tensor<?xf32> {
  %0 = "mhlo.negate"(%arg0) : (tensor<?xf32>) -> tensor<?xf32>
  func.return %0 : tensor<?xf32>
}

// CHECK-LABEL:   func @sigmoid(
// CHECK-SAME:                  %[[VAL_0:.*]]: tensor<2xf32>) -> tensor<2xf32> {
// CHECK-DAG:       %[[VAL_1:.*]] = "tf.Const"() <{value = dense<5.000000e-01> : tensor<f32>}> : () -> tensor<f32>

	}

	// SEC 1, Version 2.0, Section 4.1.4

	r, err := bigmod.NewNat().SetBytes(rBytes, c.N)
	if err != nil || r.IsZero() == 1 {
		return false
	}
	s, err := bigmod.NewNat().SetBytes(sBytes, c.N)
	if err != nil || s.IsZero() == 1 {
		return false
	}

	e := bigmod.NewNat()
	hashToNat(c, e, hash)

	// w = s⁻¹
	w := bigmod.NewNat()
	inverse(c, w, s)

	// p₁ = [e * s⁻¹]G

    FixedOutputRangeInterface,
    QuantizableResult,
    DeclareOpInterfaceMethods<TFL_ArithmeticCount>]> {
  let summary = "Logistic operator";

  let description = [{
    Computes element-wise Sigmoid of input
  }];

  let arguments = (ins TFL_TensorOf<[F32, QI8, QUI8, QI16, TFL_Quint8]>:$x);

  let results = (outs TFL_TensorOf<[F32, QI8, QUI8, QI16, TFL_Quint8]>:$y);

  let extraClassDeclaration = [{

module std/crypto/internal/bigmod/_asm

go 1.19

require github.com/mmcloughlin/avo v0.4.0

require (
	golang.org/x/mod v0.4.2 // indirect
	golang.org/x/sys v0.0.0-20211030160813-b3129d9d1021 // indirect
	golang.org/x/tools v0.1.7 // indirect
	golang.org/x/xerrors v0.0.0-20200804184101-5ec99f83aff1 // indirect

// Use of this source code is governed by a BSD-style
// license that can be found in the LICENSE file.

//go:build purego || !(386 || amd64 || arm || arm64 || ppc64 || ppc64le || riscv64 || s390x)

package bigmod

import "unsafe"

func addMulVVW1024(z, x *uint, y uint) (c uint) {
	return addMulVVW(unsafe.Slice(z, 1024/_W), unsafe.Slice(x, 1024/_W), y)
}

func addMulVVW1536(z, x *uint, y uint) (c uint) {

import (
	"strconv"

	. "github.com/mmcloughlin/avo/build"
	. "github.com/mmcloughlin/avo/operand"
	. "github.com/mmcloughlin/avo/reg"
)

//go:generate go run . -out ../nat_amd64.s -pkg bigmod

func main() {
	Package("crypto/internal/bigmod")
	ConstraintExpr("!purego")

	addMulVVW(1024)
	addMulVVW(1536)
	addMulVVW(2048)

	Generate()
}

func addMulVVW(bits int) {
	if bits%64 != 0 {