}
}

// Keep this code in sync with cmd/compile/internal/walk/builtin.go:walkUnsafeSlice
func unsafeslice64(et *_type, ptr unsafe.Pointer, len64 int64) {
	len := int(len64)
	if int64(len) != len64 {
		panicunsafeslicelen1(getcallerpc())
	}
	unsafeslice(et, ptr, len)
}

func unsafeslicecheckptr(et *_type, ptr unsafe.Pointer, len64 int64) {
	unsafeslice64(et, ptr, len64)

  tensorflow::PartialTensorShape unknown_shape;
  TensorSpec ts1;
  ts1.shape = unknown_shape;
  ts1.dtype = tensorflow::DT_FLOAT;

  TensorSpec ts2;
  ts2.shape = tensorflow::PartialTensorShape({2});
  ts2.dtype = tensorflow::DT_FLOAT;

  TensorSpec ts3;
  ts3.shape = tensorflow::PartialTensorShape({1, 2});
  ts3.dtype = tensorflow::DT_FLOAT;

  EXPECT_TRUE(absl::VerifyTypeImplementsAbslHashCorrectly({ts1, ts2, ts3}));
}

// RUN: odml-to-stablehlo-opt --compose-uniform-quantized-type \
// RUN:     --split-input-file --verify-diagnostics %s | FileCheck %s

module {
// CHECK-LABEL: quantized_conv_op
// CHECK-SAME: %[[ARG:.*]]: tensor<1x3x3x4xf32>
  func.func @quantized_conv_op(%arg0: tensor<1x3x3x4xf32>) -> tensor<1x3x3x4xf32> {
    %1 = stablehlo.constant dense<1.000000e+03> : tensor<1x1x1x1xf32>  // Input inverse scale.

  auto dtype = input->DataType();
  if (DataTypeIsFloating(BaseType(dtype)) ||
      DataTypeIsInteger(BaseType(dtype))) {
    return tensorflow::ops::Identity(ctx, input, output, name);
  } else if (!DataTypeIsComplex(BaseType(dtype)) &&
             BaseType(dtype) != DT_VARIANT) {
    return errors::InvalidArgument(
        "Expected numeric or variant tensor, got dtype ", dtype);
  }

        }
      }
    }
  }
}
node {
  name: "BoxPredictor_4/ClassPredictor/weights"
  op: "Const"
  attr {
    key: "dtype"
    value {
      type: DT_FLOAT
    }
  }
  attr {
    key: "value"
    value {
      tensor {
        dtype: DT_FLOAT
        tensor_shape {
          dim {
            size: 1
          }
          dim {
            size: 1
          }

node {
  name: "Placeholder"
  op: "Placeholder"
  attr {
    key: "dtype"
    value {
      type: DT_FLOAT
    }
  }
  attr {
    key: "shape"
    value {
      shape {
        dim {
          size: 2
        }
        dim {
          size: 5
        }
        dim {

  typedef internal::Types$i<$for j, [[T$j]]> type;
};

]]

namespace internal {

# define GTEST_TEMPLATE_ template <typename T> class

// The template "selector" struct TemplateSel<Tmpl> is used to
// represent Tmpl, which must be a class template with one type
// parameter, as a type.  TemplateSel<Tmpl>::Bind<T>::type is defined
// as the type Tmpl<T>.  This allows us to actually instantiate the

// wbMove performs the write barrier operations necessary before
// copying a region of memory from src to dst of type typ.
// Does not actually do the copying.
//
//go:nowritebarrierrec
//go:nosplit
func wbMove(typ *_type, dst, src unsafe.Pointer) {
	// This always copies a full value of type typ so it's safe to
	// pass a type here.
	//
	// See the comment on bulkBarrierPreWrite.

    AllocatorAttributes attr;
    if (cc_input.dtype() == ::tensorflow::DT_VARIANT) {
      attr.set_on_host(true);
    }
    TF_RETURN_IF_ERROR(cc_ctx->allocate_temp(cc_input.dtype(), cc_input.shape(),
                                             cc_out, attr));

    switch (cc_input.dtype()) {
      case ::tensorflow::DT_INVALID: {
        *cc_out = Tensor(::tensorflow::DT_INVALID);

  *graph.add_node() = MakeNode("a", "Placeholder", {},
                               {{"dtype", DT_FLOAT}, {"shape", input_shape}});
  *graph.add_node() = MakeNode("b", "Placeholder", {},
                               {{"dtype", DT_FLOAT}, {"shape", input_shape}});
  *graph.add_node() = MakeNode("c", "Placeholder", {},
                               {{"dtype", DT_FLOAT}, {"shape", input_shape}});