}

//===----------------------------------------------------------------------===//
// CastOp
//===----------------------------------------------------------------------===//

OpFoldResult CastOp::fold(FoldAdaptor) {
  // Cast with the same type is a no-op.
  Value operand = getOperand();
  if (getType() == operand.getType()) return operand;
  return {};

  results.add<FoldPseudoConstOp>(context);
}

//===----------------------------------------------------------------------===//
// CastOp
//===----------------------------------------------------------------------===//

OpFoldResult CastOp::fold(FoldAdaptor adaptor) {
  auto operands = adaptor.getOperands();
  assert(operands.size() == 1);
  if (getInput().getType() == getType()) {

  auto x_type = mlir::dyn_cast_or_null<ShapedType>(x.getType());
  if (!x_type) llvm_unreachable("unsupported type");
  Type type = x_type.clone(builder->getI32Type());
  return builder->create<TF::CastOp>(loc, type, x, truncate);
}
}  // namespace

//===----------------------------------------------------------------------===//
// The actual PrepareTF Pass.
//

        loc, op.getValue(), op.getBias(), feature_dim, rewriter);
    Value add = rewriter.create<AddOp>(loc, op.getValue(), bias_broadcast);
    if (add.getType() != op.getType()) {
      add = rewriter.create<tensor::CastOp>(loc, op.getType(), add);
    }
    rewriter.replaceOp(op, {add});
    return success();
  }
};

// Conterts tf.Conv2D to mhlo.dynamic_conv.

      if (body_arg.getType() == cond_arg.getType() &&
          body_arg.getType() != while_operand.getType()) {
        changed = true;
        rewriter.setInsertionPoint(while_op);
        auto cast_op = rewriter.create<CastOp>(
            while_op.getLoc(), body_arg.getType(), while_operand);
        while_op.setOperand(op_idx, cast_op);
      }
    }
    return success(changed);
  }
};

            mlir::tf_saved_model::GetBoundInputArgTypeFor(global_tensor);
        arg.setType(new_type);
        if (global_tensor.getIsMutable()) {
          auto arg_with_original_type = builder.create<mlir::TF::CastOp>(
              global_tensor.getLoc(), old_type, arg,
              /*Truncate=*/builder.getBoolAttr(false));
          arg.replaceAllUsesWith(arg_with_original_type);

	}
	switch raw.typ {
	case EvBatch:
		p.lastGs[p.lastP] = p.lastG
		if raw.args[0] != math.MaxUint64 && raw.args[0] > math.MaxInt32 {
			return fmt.Errorf("processor ID %d is larger than maximum of %d", raw.args[0], uint64(math.MaxInt32))
		}
		if raw.args[0] == math.MaxUint64 {
			p.lastP = -1
		} else {
			p.lastP = int32(raw.args[0])
		}
		p.lastG = p.lastGs[p.lastP]
		p.lastTs = Timestamp(raw.args[1])
	case EvFrequency:

	ASHA512SU1
	ASMADDL
	ASMC
	ASMNEGL
	ASMSUBL
	ASMULH
	ASMULL
	ASTLR
	ASTLRB
	ASTLRH
	ASTLRW
	ASTLXP
	ASTLXPW
	ASTLXR
	ASTLXRB
	ASTLXRH
	ASTLXRW
	ASTP
	ASTPW
	ASTXP
	ASTXPW
	ASTXR
	ASTXRB
	ASTXRH
	ASTXRW
	ASUB
	ASUBS
	ASUBSW
	ASUBW
	ASVC
	ASWPAB
	ASWPAD
	ASWPAH
	ASWPALB
	ASWPALD
	ASWPALH
	ASWPALW
	ASWPAW

	{ASTP, C_PAIR, C_NONE, C_NONE, C_NPOREG, C_NONE, 67, 4, 0, 0, 0},
	{ASTP, C_PAIR, C_NONE, C_NONE, C_NPOREG, C_NONE, 67, 4, 0, 0, C_XPRE},
	{ASTP, C_PAIR, C_NONE, C_NONE, C_NPOREG, C_NONE, 67, 4, 0, 0, C_XPOST},
	{ASTP, C_PAIR, C_NONE, C_NONE, C_PPOREG, C_NONE, 67, 4, 0, 0, 0},
	{ASTP, C_PAIR, C_NONE, C_NONE, C_PPOREG, C_NONE, 67, 4, 0, 0, C_XPRE},
	{ASTP, C_PAIR, C_NONE, C_NONE, C_PPOREG, C_NONE, 67, 4, 0, 0, C_XPOST},

				q1.Reg = REGSP
				q1.To.Type = obj.TYPE_REG
				q1.To.Reg = REG_R20

				prologueEnd = q1

				// STP (R29, R30), -8(R20)
				q1 = obj.Appendp(q1, c.newprog)
				q1.Pos = p.Pos
				q1.As = ASTP
				q1.From.Type = obj.TYPE_REGREG
				q1.From.Reg = REGFP
				q1.From.Offset = REGLINK
				q1.To.Type = obj.TYPE_MEM
				q1.To.Reg = REG_R20
				q1.To.Offset = -8