// Create new ConstantOp-ConvertOp-Operation sequences. At this moment,
      // old ConstantOp is guaranteed to have one F32->F16 convert op regardless
      // of its number of users.
      rewriter.setInsertionPointAfter(op);
      // create new F16 constant op in that location
      ConstantOp new_const = rewriter.create<ConstantOp>(
          op->getLoc(), new_result_type, new_value_attr);

	{"V23", "V23"},
	{"V24", "V24"},
	{"V25", "V25"},
	{"V26", "V26"},
	{"V27", "V27"},
	{"V28", "V28"},
	{"V29", "V29"},
	{"V30", "V30"},
	{"V31", "V31"},
	{"F14", "F14"},
	{"F15", "F15"},
	{"F16", "F16"},
	{"F17", "F17"},
	{"F18", "F18"},
	{"F19", "F19"},
	{"F20", "F20"},
	{"F21", "F21"},
	{"F22", "F22"},
	{"F23", "F23"},
	{"F24", "F24"},
	{"F25", "F25"},
	{"F26", "F26"},

	FABSS F2, F28                              // 5cc0201e
	FABSD F0, F14                              // 0ec0601e
	//TODO FACGE F25, F16, F0                  // 00ee797e
	//TODO VFACGE V11.S2, V15.S2, V9.S2        // e9ed2b2e
	//TODO FACGT F20, F16, F27                 // 1beef47e
	//TODO VFACGT V15.S4, V25.S4, V22.S4       // 36efaf6e
	//TODO VFADD V21.D2, V10.D2, V21.D2        // 55d5754e

  %5 = "tfl.mul"(%0, %1) {fused_activation_function = "NONE"} : (tensor<  f16>, tensor<  f16>) -> tensor<  f16>
  %6 = "tfl.mul"(%0, %3) {fused_activation_function = "NONE"} : (tensor<  f16>, tensor<4xf16>) -> tensor<4xf16>
  %7 = "tfl.mul"(%2, %1) {fused_activation_function = "NONE"} : (tensor<4xf16>, tensor<  f16>) -> tensor<4xf16>

  func.return %0 : tensor<4xcomplex<f64>>
}

func.func @f16() -> tensor<4xf16> {
  // CHECK-LABEL: @f16
  // CHECK: value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf16>
  %0 = "tfl.pseudo_const"() { value = dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf16> } : () -> tensor<4xf16>

  func.return %0 : tensor<4xcomplex<f64>>
}

func.func @f16() -> tensor<4xf16> {
  // CHECK-LABEL: @f16
  // CHECK: value = dense<[1.000000e+00, 2.000000e+00, 3.000000e+00, 4.000000e+00]> : tensor<4xf16>
  %0 = "tfl.pseudo_const"() { value = dense<[1.0, 2.0, 3.0, 4.0]> : tensor<4xf16> } : () -> tensor<4xf16>

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

	"R12",
	"R13",
	"R14",
	"R15",

	"F0",
	"F1",
	"F2",
	"F3",
	"F4",
	"F5",
	"F6",
	"F7",
	"F8",
	"F9",
	"F10",
	"F11",
	"F12",
	"F13",
	"F14",
	"F15",

	"F16",
	"F17",
	"F18",
	"F19",
	"F20",
	"F21",
	"F22",
	"F23",
	"F24",
	"F25",
	"F26",
	"F27",
	"F28",
	"F29",
	"F30",
	"F31",

	"SP",
	"g",

	// pseudo-registers

        tensor<256xf64>) -> tensor<4x256xf64>
  func.return %0 : tensor<4x256xf64>
}

// CHECK-LABEL: @batchNormInference_f16
// Validate that epsilon is properly down to f16
// CHECK-DAG: %[[EPS:.+]] = mhlo.constant dense<1.000000e+00> : tensor<256xf16>
func.func @batchNormInference_f16(
    %x: tensor<4x256xf16>, %scale: tensor<256xf16>, %offset: tensor<256xf16>,

		// copy should fail.
		name: "barrierslice",
		c: `#include <stdlib.h>
		    struct s16 { char *a[10]; };
		    struct s16 *f16() { return malloc(sizeof(struct s16)); }
		    void f16b(struct s16 *p) {}`,
		body:      `p := C.f16(); copy(p.a[:], []*C.char{new(C.char)}); C.f16b(p)`,
		fail:      true,
		expensive: true,
	},
	{
		// A very large value uses a GC program, which is a