// -----

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

  // borrowing seats from this level.
  // The server's concurrency limit (ServerCL) is divided among the
  // Limited priority levels in proportion to their NCS values:
  //
  // NominalCL(i)  = ceil( ServerCL * NCS(i) / sum_ncs )
  // sum_ncs = sum[limited priority level k] NCS(k)
  //
  // Bigger numbers mean a larger nominal concurrency limit,
  // at the expense of every other Limited priority level.

}

// alignDown rounds n down to a multiple of a. a must be a power of 2.
//
//go:nosplit
func alignDown(n, a uintptr) uintptr {
	return n &^ (a - 1)
}

// divRoundUp returns ceil(n / a).
func divRoundUp(n, a uintptr) uintptr {
	// a is generally a power of two. This will get inlined and
	// the compiler will optimize the division.
	return (n + a - 1) / a
}

      return logical_device;
    }

    if (  // Cast, real/imag, etc.
        def->hasTrait<
            mlir::OpTrait::TF::SameOperandsAndResultTypeResolveRef>() ||
        // Exp, ceil, etc.
        def->hasTrait<mlir::OpTrait::SameOperandsAndResultType>() ||
        // Identity
        def->hasTrait<mlir::OpTrait::TF::OperandsSameAsResultsTypeOrRef>() ||
        // AddV2, Sub, etc.

		// will be rounded up which most likely will lead to the last subset having less
		// endpoints than the expected proportion.
		toBeAdded := int(math.Ceil((float64(numInSubset) / float64(numTotal)) * float64(max)))
		// If there is not enough endpoints for the last subset, ensure only the number up
		// to the capacity are added
		if toBeAdded > canBeAdded {

  FILL = 94,
  FLOOR_MOD = 95,
  RANGE = 96,
  RESIZE_NEAREST_NEIGHBOR = 97,
  LEAKY_RELU = 98,
  SQUARED_DIFFERENCE = 99,
  MIRROR_PAD = 100,
  ABS = 101,
  SPLIT_V = 102,
  UNIQUE = 103,
  CEIL = 104,
  REVERSE_V2 = 105,
  ADD_N = 106,
  GATHER_ND = 107,
  COS = 108,
  WHERE = 109,
  RANK = 110,
  ELU = 111,
  REVERSE_SEQUENCE = 112,
  MATRIX_DIAG = 113,
  QUANTIZE = 114,

	// We don't have an exact data, but given we store updates from
	// the last <eventFreshDuration>, we approach it by dividing the
	// capacity by the length of the history window.
	chanSize := int(math.Ceil(float64(w.currentCapacity()) / eventFreshDuration.Seconds()))

	// Finally we adjust the size to avoid ending with too low or
	// to large values.
	if chanSize < minWatchChanSize {
		chanSize = minWatchChanSize
	}

    int num_bits = qtype.getStorageTypeIntegralWidth();
    if (num_bits == 8) {
      // If storage is 8-bit, trained num bits may be less than 8 so check here.
      num_bits =
          static_cast<int>(std::ceil(std::log2(qtype.getStorageTypeMax())));
    }
    // This is a positive value, and will be applied on zero points and fixed
    // point ranges.
    int64_t offset =