VPMSUMD	V1,const1,V1
	VAND	V1,mask_32bit,V1
	VPMSUMD	V1,const2,V1
	VXOR	V0,V1,V0
	VSLDOI  $4,V0,zeroes,V0

#endif

	MFVSRD	VS32,R3 // VS32 = V0

	NOR	R3,R3,R3 // return ^crc
	MOVW	R3,ret+32(FP)
	RET

first_warm_up_done:

	LVX	(R3),const1
	ADD	$16,R3

	VPMSUMD	V16,const1,V8
	VPMSUMD	V17,const1,V9
	VPMSUMD	V18,const1,V10
	VPMSUMD	V19,const1,V11
	VPMSUMD	V20,const1,V12

	// First reduction step
	MUL	acc0, hlp1, hlp0

	MUL	const0, hlp1, t0
	ADDS	t0, acc0, acc0
	UMULH	const0, hlp0, t1

	MUL	const1, hlp0, t0
	ADCS	t0, acc1, acc1
	UMULH	const1, hlp0, y0

	MUL	const2, hlp0, t0
	ADCS	t0, acc2, acc2
	UMULH	const2, hlp0, acc0

	MUL	const3, hlp0, t0
	ADCS	t0, acc3, acc3

	UMULH	const3, hlp0, hlp0
	ADC	$0, hlp0

	ADDS	t1, acc1, acc1

  // CHECK-DAG: [[cst1:%.*]] = arith.constant dense<-16> : tensor<4xi32>
  // CHECK-DAG: [[cst2:%.*]] = arith.constant dense<4> : tensor<4xi32>
  // CHECK-DAG: [[cst3:%.*]] = arith.constant dense<-8> : tensor<4xi32>
  // CHECK: return [[cst0]], [[cst1]], [[cst2]], [[cst3]]

  %5 = "tfl.mul"(%0, %1) {fused_activation_function = "NONE"} : (tensor<  i32>, tensor<  i32>) -> tensor<  i32>

		var cost uint64
		if len(args) > 0 {
			cost += traversalCost(args[0]) // these O(n) operations all cost roughly the cost of a single traversal
		}
		return &cost
	case "url", "lowerAscii", "upperAscii", "substring", "trim":
		if len(args) >= 1 {
			cost := uint64(math.Ceil(float64(actualSize(args[0])) * common.StringTraversalCostFactor))
			return &cost
		}
	case "replace", "split":
		if len(args) >= 1 {

// CHECK-DAG: %[[CONST_0:.*]] = "tf.Const"() <{value = dense<{{.*}}> : tensor<1x1x3x2xf32>}> : () -> tensor<1x1x3x2xf32>
// CHECK-DAG: %[[CONST_1:.*]] = "tf.Const"() <{value = dense<-1.000000e+00> : tensor<f32>}> : () -> tensor<f32>
// CHECK-DAG: %[[CONST_2:.*]] = "tf.Const"() <{value = dense<1.000000e+00> : tensor<f32>}> : () -> tensor<f32>

absl::StatusOr<Node*> ReplaceFunctionCallWithPartitionedCall(
    const GraphOptimizationPassOptions& options,
    const FunctionLibraryDefinition& flib_def, Node* n, Graph* g,
    const NameAttrList& func, const Scope& root) {
  string config_string = options.session_options->config.SerializeAsString();

  int input_count = absl::c_count_if(
      n->in_edges(), [](const Edge* e) { return !e->IsControlEdge(); });

// COM: CHECK-DAG: %[[CONST:.*]] = stablehlo.constant dense<{{.*}}> : tensor<1x1x2x8xf32>
// COM: CHECK: return %[[CONST]] : tensor<1x1x2x8xf32>

func.func @fused_batchnorm_no_training_arg_input(%arg_0: tensor<1x1x2x8xf32>) -> (tensor<1x1x2x8xf32>) {
  %cst_0 = "tf.Const"() {value = dense<[0.1, 0.2, 0.1, 0.2, 0.1, 0.2, 0.1, 0.2]> : tensor<8xf32>} : () -> tensor<8xf32>

class QuantizeConstWeights : public OpRewritePattern<TF::ConstOp> {
 public:
  explicit QuantizeConstWeights(
      MLIRContext* context,
      const tensorflow::quantization::QuantizationOptions& quantization_options)
      : OpRewritePattern<TF::ConstOp>(context),
        quant_options_(quantization_options) {}

  LogicalResult matchAndRewrite(TF::ConstOp op,

// Creates a pass that creates a RestoreV2 op in the initializer function with
// type "restore_op" that initializes variables from the checkpoint. It finds
// tf.AssignVariableOp(tf.VarHandleOp, tf.Const) patterns in the initializer
// function and replaces tf.Consts with the results of RestoreV2.
std::unique_ptr<OperationPass<ModuleOp>> CreateInsertRestoreOpPass();

// Creates a pass that creates a new function that wraps the newly created

StatusOrOptional<SliceInputs> GetSliceInputs(Node* slice) {
  const int kSliceInputIndex = 0;
  const int kSliceBeginIndex = 1;
  const int kSliceSizeIndex = 2;

  const Edge* slice_input_edge;
  TF_RETURN_IF_ERROR(slice->input_edge(kSliceInputIndex, &slice_input_edge));
  const Edge* slice_size_edge;
  TF_RETURN_IF_ERROR(slice->input_edge(kSliceSizeIndex, &slice_size_edge));
  const Edge* slice_begin_edge;