P3 = -1.228866684490136173410e+02
		P4 = -6.485021904942025371773e+01
		Q0 = +2.485846490142306297962e+01
		Q1 = +1.650270098316988542046e+02
		Q2 = +4.328810604912902668951e+02
		Q3 = +4.853903996359136964868e+02
		Q4 = +1.945506571482613964425e+02
	)
	z := x * x
	z = z * ((((P0*z+P1)*z+P2)*z+P3)*z + P4) / (((((z+Q0)*z+Q1)*z+Q2)*z+Q3)*z + Q4)
	z = x*z + x
	return z
}

			// When Q(k) > 1e9	good for double
			// When Q(k) > 1e17	good for quadruple

			// determine k
			w := float64(n+n) / x
			h := 2 / x
			q0 := w
			z := w + h
			q1 := w*z - 1
			k := 1
			for q1 < 1e9 {
				k++
				z += h
				q0, q1 = q1, z*q1-q0
			}
			m := n + n
			t := 0.0
			for i := 2 * (n + k); i >= m; i -= 2 {
				t = 1 / (float64(i)/x - t)
			}
			a := t
			b = 1

		rhat += yn1
		if rhat >= two32 {
			break
		}
	}

	un21 := un32*two32 + un1 - q1*y
	q0 := un21 / yn1
	rhat = un21 - q0*yn1

	for q0 >= two32 || q0*yn0 > two32*rhat+un0 {
		q0--
		rhat += yn1
		if rhat >= two32 {
			break
		}
	}

	return q1*two32 + q0, (un21*two32 + un0 - q0*y) >> s
}

// Rem returns the remainder of (hi, lo) divided by y. Rem panics for

//         for x in (0,2)
//              j0(x) = 1-z/4+ z**2*R0/S0,  where z = x*x;
//         (precision:  |j0-1+z/4-z**2R0/S0 |<2**-63.67 )
//         for x in (2,inf)
//              j0(x) = sqrt(2/(pi*x))*(p0(x)*cos(x0)-q0(x)*sin(x0))
//         where x0 = x-pi/4. It is better to compute sin(x0),cos(x0)
//         as follow:
//              cos(x0) = cos(x)cos(pi/4)+sin(x)sin(pi/4)
//                      = 1/sqrt(2) * (cos(x) + sin(x))

			op = "FMOVS" + suffix
			args[0] = fmt.Sprintf("F%d", rno&31)
		} else if rno >= uint16(D0) && rno <= uint16(D31) {
			op = "FMOVD" + suffix
			args[0] = fmt.Sprintf("F%d", rno&31)
		} else if rno >= uint16(Q0) && rno <= uint16(Q31) {
			op = "FMOVQ" + suffix
			args[0] = fmt.Sprintf("F%d", rno&31)
		} else {
			op = "MOVD" + suffix
		}

	case LDRB:
		op = "MOVBU" + suffix

	case LDRH:

// CHECK: %[[Q0:.*]] = "tfl.quantize"(%arg0) <{qtype = tensor<1x2x!quant.uniform<u8:f32, 2.000000e+00:128>>}> : (tensor<1x2xf32>) -> tensor<1x2x!quant.uniform<u8:f32, 2.000000e+00:128>>

		return fmt.Sprintf("H%d", int(r-H0))
	case S0 <= r && r <= S31:
		return fmt.Sprintf("S%d", int(r-S0))
	case D0 <= r && r <= D31:
		return fmt.Sprintf("D%d", int(r-D0))
	case Q0 <= r && r <= Q31:
		return fmt.Sprintf("Q%d", int(r-Q0))

	case V0 <= r && r <= V31:
		return fmt.Sprintf("V%d", int(r-V0))
	default:
		return fmt.Sprintf("Reg(%d)", int(r))
	}
}

	inst *syntax.Inst
	cap  []int
}

// A machine holds all the state during an NFA simulation for p.
type machine struct {
	re       *Regexp      // corresponding Regexp
	p        *syntax.Prog // compiled program
	q0, q1   queue        // two queues for runq, nextq
	pool     []*thread    // pool of available threads
	matched  bool         // whether a match was found
	matchcap []int        // capture information for the match

  // CHECK-DAG: %[[cst2:.*]] = stablehlo.constant dense<6.000000e+00> : tensor<f32>
  // CHECK-DAG: %[[cst3:.*]] = stablehlo.constant dense<0.000000e+00> : tensor<f32>
  // CHECK-NOT: %[[q0:.*]] = "quantfork.qcast"(%[[cst0]])
  // CHECK-NOT: %[[q1:.*]] = "quantfork.qcast"(%[[cst1]])
  // CHECK: %[[q2:.*]] = "quantfork.qcast"(%[[cst2]])
  // CHECK-SAME: quant.uniform<i8:f32, 0.023529411764705882:-128>

  func.return %1 : tensor<2x2x!quant.uniform<u8:f32, 1.000000e-01:128>>

// CHECK: %[[q0:.*]] = "tfl.quantize"(%arg1) <{qtype = tensor<1x2x!quant.uniform<u8:f32, 1.000000e-01:128>>}> {volatile}
// CHECK: %[[q1:.*]] = "tfl.quantize"(%arg0) <{qtype = tensor<1x2x!quant.uniform<u8:f32, 1.000000e-01:128>>}> {volatile}
// CHECK: %[[cc:.*]] = "tfl.concatenation"(%[[q1]], %[[q0]]) <{axis = 0 : i32, fused_activation_function = "NONE"}>