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

}

type T1 struct {
	X *int
}

func f15(x []T1, y T1) []T1 {
	return append(x, y) // ERROR "write barrier"
}

type T8 struct {
	X [8]*int
}

func f16(x []T8, y T8) []T8 {
	return append(x, y) // ERROR "write barrier"
}

func t1(i interface{}) **int {
	// From issue 14306, make sure we have write barriers in a type switch
	// where the assigned variable escapes.

00000000  14 03 03 00 01 01 16 03  03 00 20 2e d7 9e b0 66  |.......... ....f|
00000010  f1 39 84 a9 d9 93 17 bd  94 5a 31 49 bb eb f0 a0  |.9.......Z1I....|
00000020  7b af d0 3b ae 1a 5d f6  46 31 36                 |{..;..].F16|
>>> Flow 5 (client to server)
00000000  17 03 03 00 16 cd 39 a7  64 3b 6a de 14 e0 26 ea  |......9.d;j...&.|
00000010  66 b2 73 b1 8e b0 e3 a9  94 62 4f                 |f.s......bO|
>>> Flow 6 (server to client)

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

    return true;
  };
  return just_call(while_op.getBody()) && just_call(while_op.getCond());
}

bool IsCompatibleTypeWithTFLCastOp(Type type) {
  auto elemType = getElementTypeOrSelf(type);
  // F16, F32, F64, BF16 types are allowed.
  if (elemType.isBF16() || elemType.isF16() || elemType.isF32() ||
      elemType.isF64())
    return true;

  // I1, I4, I8, I16, I32, I64 types are allowed.

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

	// R30 is REGTMP not used in regalloc
	"R31",

	"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",

	"g",   // aka R30
	"R31", // aka REGLINK

	"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",

	"HI", // high bits of multiplication

	"R28",
	"SP",  // aka R29
	"g",   // aka R30
	"R31", // REGLINK

	// odd FP registers contain high parts of 64-bit FP values
	"F0",
	"F2",
	"F4",
	"F6",
	"F8",
	"F10",
	"F12",
	"F14",
	"F16",
	"F18",
	"F20",
	"F22",
	"F24",
	"F26",
	"F28",
	"F30",

	"HI", // high bits of multiplication
	"LO", // low bits of multiplication

	// If you add registers, update asyncPreempt in runtime.