dummy()
	}

	// s390x:"CIJ\t[$]4, R[0-9]+, [$]-128, "
	for i := doNotOptimize; i <= -129; i-- {
		dummy()
	}
}

// Unsigned 32-bit compare-and-branch.
func ui32(x, y chan uint32) {
	// s390x:"CLRJ\t[$](2|4), R[0-9]+, R[0-9]+, "
	for <-x > <-y {
		dummy()
	}

	// s390x:"CL?RJ\t[$]6, R[0-9]+, R[0-9]+, "
	for <-x != <-y {
		dummy()
	}
}

func TestEmptyMapWithInterfaceKey(t *testing.T) {
	var (
		b    bool
		i    int
		i8   int8
		i16  int16
		i32  int32
		i64  int64
		ui   uint
		ui8  uint8
		ui16 uint16
		ui32 uint32
		ui64 uint64
		uipt uintptr
		f32  float32
		f64  float64
		c64  complex64
		c128 complex128
		a    [4]string
		s    string
		p    *int
		up   unsafe.Pointer

include "stablehlo/dialect/ChloOps.td"
include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.td"
include "mhlo/IR/hlo_ops.td"

def SignedIntTensor : TensorOf<[I1, I8, I16, I32, I64]>;
def UnsignedIntTensor : TensorOf<[UI8, UI16, UI32, UI64]>;

// IEEE compliant floating point tensors.
def IEEEFloatTensor : TensorOf<[F16, F32, F64]>;

//===----------------------------------------------------------------------===//
// BatchNorm op patterns.

// CHECK:         }
func.func @convert_dynamic_slice_ui32(%arg0: tensor<7x3xf32>, %arg1: tensor<ui32>, %arg2: tensor<ui32>) -> tensor<4x2xf32> {
  %0 = "mhlo.dynamic_slice"(%arg0, %arg1, %arg2) <{slice_sizes = dense<[4, 2]> : tensor<2xi64>}> : (tensor<7x3xf32>, tensor<ui32>, tensor<ui32>) -> tensor<4x2xf32>
  func.return %0 : tensor<4x2xf32>
}

// CHECK-LABEL:   func.func @convert_scatter_update(

}

// CHECK-LABEL: testMul32BitUInt
func.func @testMul32BitUInt(tensor<? x ui32>, tensor<? x ui32>) -> tensor<? x ui32> {
^bb0(%arg0: tensor<? x ui32>, %arg1: tensor<? x ui32>):
  // CHECK: tfl.mul %arg0, %arg1 {fused_activation_function = "RELU6"}
  %0 = tfl.mul %arg0, %arg1 {fused_activation_function = "RELU6"} : tensor<? x ui32>
  func.return %0#0 : tensor<? x ui32>
}

// CHECK-LABEL: testMul16BitInt

func.func @relu6_unsigned(%arg0: tensor<?xui32>) -> tensor<?xui32> {
  // CHECK-DAG: %[[ZERO:.*]] = mhlo.constant dense<0> : tensor<ui32>
  // CHECK-DAG: %[[SIX:.*]] = mhlo.constant dense<6> : tensor<ui32>
  // CHECK: mhlo.clamp %[[ZERO]], %arg0, %[[SIX]] : (tensor<ui32>, tensor<?xui32>, tensor<ui32>) -> tensor<?xui32>
  %0 = "tf.Relu6"(%arg0) : (tensor<?xui32>) -> tensor<?xui32>
  func.return %0: tensor<?xui32>
}

// -----

  TF_DerivedResultTypeAttr dtype = TF_DerivedResultTypeAttr<0>;
}

// TODO(lyandy): Investigate supported dtypes (`minval`, `maxval`, `output`) for
// `tf.StatefulUniformInt`. tf2xla kernels support i32, i64, ui32, and ui64
// while TensorFlow CPU/GPU kernels only support i32 and i64.
def TF_StatefulUniformIntOp : TF_Op<"StatefulUniformInt", []> {
  let summary = "Outputs random integers from a uniform distribution.";

func.func @broadcast_to_ui32(%arg0: tensor<ui32>, %arg1: tensor<1xi64>) -> tensor<10xui32> {
  %0 = "tfl.broadcast_to"(%arg0, %arg1) : (tensor<ui32>, tensor<1xi64>) -> tensor<10xui32>
  return %0 : tensor<10xui32>
  // CHECK:  %cst = arith.constant dense<1> : tensor<10xui32>
  // CHECK:  %0 = tfl.mul(%arg0, %cst) <{fused_activation_function = "NONE"}> : (tensor<ui32>, tensor<10xui32>) -> tensor<10xui32>

}

// Return true when the given element_type is I32.
bool IsI32Type(Type element_type) {
  return element_type.isInteger(32) && !element_type.isUnsignedInteger();
}

// Return true when the given element_type is UI32.
bool IsUI32Type(Type element_type) {
  return element_type.isInteger(32) && element_type.isUnsignedInteger();
}

// Return true when the given element_type is I64.
bool IsI64Type(Type element_type) {