//	...
//	wait.BackoffUntil(..., bm.Backoff, ...)
//
// Use:
//
//	delayFn := wait.Backoff{
//	  Duration: init,
//	  Cap:      max,
//	  Steps:    int(math.Ceil(float64(max) / float64(init))), // now a required argument
//	  Factor:   factor,
//	  Jitter:   jitter,
//	}.DelayWithReset(reset, clock)
//	wait.BackoffUntil(..., delayFn.Timer(), ...)

(LocalAddr <t> {sym} base _)  && !t.Elem().HasPointers() => (LoweredAddr {sym} base)

// Write barrier.
(WB ...) => (LoweredWB ...)

// --- Intrinsics ---
(Sqrt ...) => (F64Sqrt ...)
(Trunc ...) => (F64Trunc ...)
(Ceil ...) => (F64Ceil ...)
(Floor ...) => (F64Floor ...)
(RoundToEven ...) => (F64Nearest ...)
(Abs ...) => (F64Abs ...)
(Copysign ...) => (F64Copysign ...)

(Sqrt32 ...) => (F32Sqrt ...)

		}

		// Take the ceiling of the reservation, since fractional CPUs cannot be
		// exclusively allocated.
		reservedCPUsFloat := float64(reservedCPUs.MilliValue()) / 1000
		numReservedCPUs := int(math.Ceil(reservedCPUsFloat))
		policy, err = NewStaticPolicy(topo, numReservedCPUs, specificCPUs, affinity, cpuPolicyOptions)
		if err != nil {
			return nil, fmt.Errorf("new static policy error: %w", err)
		}

	default:

	}
	trimmed := 0
	// Remember last trimmed digit to check for round-up.
	// c0 will be used to remember zeroness of following digits.
	cNextDigit := 0
	for upper > 0 {
		// Repeatedly compute:
		// l = Ceil(lower / 10^k)
		// c = Round(central / 10^k)
		// u = Floor(upper / 10^k)
		// and stop when c goes out of the (l, u) interval.
		l := (lower + 9) / 10
		c, cdigit := central/10, central%10
		u := upper / 10

    // We have q=100, k=index, and N=9951. Therefore k*(N-1)/q is 99.5*index. If index is even, that
    // is an integer 199*index/2. If index is odd, that is halfway between floor(199*index/2) and
    // ceil(199*index/2).
    if (index % 2 == 0) {
      int position = IntMath.divide(199 * index, 2, UNNECESSARY);
      return PSEUDORANDOM_DATASET_SORTED.get(position);
    } else {

  // the assured concurrency value (ACV) --- the number of requests
  // that may be executing at a time --- for each such priority
  // level:
  //
  //             ACV(l) = ceil( SCL * ACS(l) / ( sum[priority levels k] ACS(k) ) )
  //
  // bigger numbers of ACS mean more reserved concurrent requests (at the
  // expense of every other PL).
  // This field has a default value of 30.
  // +optional

// -----

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

// -----

// CHECK-LABEL: func @ceil_dynamic
func.func @ceil_dynamic(%arg0: tensor<?xf32>) -> tensor<?xf32> {
  // CHECK:  mhlo.ceil %arg0 : tensor<?xf32>

  // the assured concurrency value (ACV) --- the number of requests
  // that may be executing at a time --- for each such priority
  // level:
  //
  //             ACV(l) = ceil( SCL * ACS(l) / ( sum[priority levels k] ACS(k) ) )
  //
  // bigger numbers of ACS mean more reserved concurrent requests (at the
  // expense of every other PL).
  // This field has a default value of 30.
  // +optional

  // the assured concurrency value (ACV) --- the number of requests
  // that may be executing at a time --- for each such priority
  // level:
  //
  //             ACV(l) = ceil( SCL * ACS(l) / ( sum[priority levels k] ACS(k) ) )
  //
  // bigger numbers of ACS mean more reserved concurrent requests (at the
  // expense of every other PL).
  // This field has a default value of 30.
  // +optional

// Radzik, editors, Algorithms - ESA 2004, volume 3221 of Lecture Notes in
// Computer Science, pages 714-723. Springer, 2004.
//
// Let M = m-a and N = b-n. Wolog M < N.
// The recursion depth is bound by ceil(log(N+M)).
// The algorithm needs O(M*log(N/M + 1)) calls to data.Less.
// The algorithm needs O((M+N)*log(M)) calls to data.Swap.
//
// The paper gives O((M+N)*log(M)) as the number of assignments assuming a