// CHECK: %[[GROUP_SIZE:.*]] = mhlo.constant dense<1.000000e+00>
  // CHECK: %[[REDUCE:.*]] = "mhlo.all_reduce"
  // CHECK: mhlo.add
  // CHECK: %[[RESULT:.*]] = mhlo.divide %[[REDUCE]], %[[GROUP_SIZE]]
  // CHECK-NEXT: return %[[RESULT]]
  %0 = "tf.XlaAllReduce"(%input, %group_assignment) {reduce_op = "Mean", mode = "CrossReplica"} : (tensor<f32>, tensor<2x1xi32>) -> tensor<f32>

// RUN: tf-tfrt-opt -tfrt-lower-cluster-to-runtime-ops-tpu -split-input-file -verify-diagnostics %s | FileCheck %s

module attributes {tf.versions = {producer = 888 : i32}, tf.devices = ["/job:worker/replica:0/task:0/device:CPU:0", "/job:worker/replica:0/task:0/device:TPU_SYSTEM:0", "/job:worker/replica:0/task:0/device:TPU:0"]} {

  // CHECK-LABEL: @converts_cluster
  func.func @converts_cluster() {
    // CHECK: %0:2 = "tf_device.launch"() <{{.*}}> ({

	// rounding up (see FIPS 203 (DRAFT), Section 2.3).

	dividend := uint32(y) * q
	quotient := dividend >> d // (y * q) / 2ᵈ

	// The d'th least-significant bit of the dividend (the most significant bit
	// of the remainder) is 1 for the top half of the values that divide to the
	// same quotient, which are the ones that round up.
	quotient += dividend >> (d - 1) & 1

	return portalHostMap, nil
}

// FindDevicesForISCSILun unsupported returns nil
func (handler *deviceHandler) FindDevicesForISCSILun(targetIqn string, lun int) ([]string, error) {
	devices := []string{}
	return devices, nil

    new RoundToDoubleTester(BigDecimal.ZERO).setExpectation(0.0, values()).test();
  }

  public void testRoundToDouble_oneThird() {
    new RoundToDoubleTester(
            BigDecimal.ONE.divide(BigDecimal.valueOf(3), new MathContext(50, HALF_EVEN)))
        .roundUnnecessaryShouldThrow()
        .setExpectation(0.33333333333333337, UP, CEILING)

	return 1<<63 - 1 + int64(unsafe.Sizeof(*new(T)))
}

func main() {
	shift[[62]byte]()
	shift[[63]byte]()
	shift[[64]byte]()
	shift[[100]byte]()
	shift[[1e6]byte]()

	add[[1]byte]()
	shouldPanic("divide by zero", func() { div[[0]byte]() })
}

func shouldPanic(str string, f func()) {
	defer func() {
		err := recover()
		if err == nil {
			panic("did not panic")
		}
		s := err.(error).Error()

// CHECK-SAME:    sharding={
// CHECK-SAME:    {devices=[1,2]0,1}
// CHECK-SAME:    {maximal device=0}
// CHECK-SAME:    {replicated}
// CHECK-SAME:    }
// CHECK:         get-tuple-element((f32[128,10]{1,0}, f32[10,1024]{1,0}, f32[128,1024]{1,0}) %[[ARG_TUPLE]]), index=0
// CHECK-SAME:    sharding={devices=[1,2]0,1}

// The id of the compiled cluster.
extern const char* const kXlaClusterIdAttr;  // "_xla_compile_id"

[[deprecated("XLA:CPU/GPU devices are deprecated")]] void
RequestXlaDevicesCreation();

[[deprecated("XLA:CPU/GPU devices are deprecated")]] bool
XlaDevicesCreationRequired();

}  // namespace tensorflow

  auto& side_effect_analysis = getAnalysis<mlir::TF::SideEffectAnalysis>();
  // Get runtime devices information from the closest parent module.
  auto module = getOperation();
  mlir::TF::RuntimeDevices devices;
  if (failed(tensorflow::GetDevicesFromOp(module, &devices)))
    return signalPassFailure();

  OpBuilder builder(&getContext());
  llvm::SmallVector<mlir::tf_device::ClusterOp, 4> clusters;

const DefaultPhysPageSize = goarch.DefaultPhysPageSize

// PCQuantum is the minimal unit for a program counter (1 on x86, 4 on most other systems).
// The various PC tables record PC deltas pre-divided by PCQuantum.
const PCQuantum = goarch.PCQuantum

// Int64Align is the required alignment for a 64-bit integer (4 on 32-bit systems, 8 on 64-bit).
const Int64Align = goarch.PtrSize