//
//go:systemstack
func checkRanks(gp *g, prevRank, rank lockRank) {
	rankOK := false
	if rank < prevRank {
		// If rank < prevRank, then we definitely have a rank error
		rankOK = false
	} else if rank == lockRankLeafRank {
		// If new lock is a leaf lock, then the preceding lock can
		// be anything except another leaf lock.
		rankOK = prevRank < lockRankLeafRank
	} else {
		// We've now verified the total lock ranking, but we

}

func (rank lockRank) String() string {
	if rank == 0 {
		return "UNKNOWN"
	}
	if rank == lockRankLeafRank {
		return "LEAF"
	}
	if rank < 0 || int(rank) >= len(lockNames) {
		return "BAD RANK"
	}
	return lockNames[rank]
}

// lockPartialOrder is the transitive closure of the lock rank graph.
// An entry for rank X lists all of the ranks that can already be held

//
// $iota = "mhlo.iota"(){iota_dimension = $dimensions[0]},
//
// where $dimensions must have size 1 and iota can have rank>=1.
// It usually used for matching rank 1 iota since the iotaOp will be folded to
// IotaOp + BroadCastInDimOp except for the case when result shape is rank 1.
bool MatchSingleIota(DenseIntElementsAttr dimensions, Value iota) {
  auto iota_op = dyn_cast_or_null<mhlo::IotaOp>(iota.getDefiningOp());

  RankedTensorType tensorType = mlir::cast<RankedTensorType>(value.getType());
  Type element_type = tensorType.getElementType();

  int rank = tensorType.getShape().size();
  int num_rows = tensorType.getShape()[rank - 2];
  int num_cols = tensorType.getShape()[rank - 1];

  std::vector<Value> sliced;

  if (batch_size == 1) {
    // Batch size is 1, no splitting is required

  // Operand must have statically known rank.
  kRank = 0,
  // Operand must have statically known shape (all dimensions are known at
  // compile time).
  kShape = 1,
  // Operand must have statically known value (operand must be defined by a
  // constant operation).
  kValue = 2,
};

// Returns the more restrictive constraint of `a` and `b`:
//
//    Value >> Shape >> Rank
//

	// If one of the two pods is on the same node as one or more additional
	// ready pods that belong to the same replicaset, whichever pod has more
	// colocated ready pods is less
	if s.Rank[i] != s.Rank[j] {
		return s.Rank[i] > s.Rank[j]
	}
	// TODO: take availability into account when we push minReadySeconds information from deployment into pods,
	//       see https://github.com/kubernetes/kubernetes/issues/22065

  REGISTER_KERNEL_BUILDER(                                                     \
      Name("Rank").Device(DEVICE).HostMemory("output").TypeConstraint("T",     \
                                                                      TYPES),  \
      RankOp);                                                                 \
  REGISTER_KERNEL_BUILDER(                                                     \

  auto str = [](ValueConstraint constraint) -> StringRef {
    switch (constraint) {
      case ValueConstraint::kRank:
        return "rank";
      case ValueConstraint::kShape:
        return "shape";
      case ValueConstraint::kValue:
        return "value";
      default:
        llvm_unreachable("unknown value constraint");
    }
  };

  os << str(constraint);
  return os;
}

    SmallVector<int64_t> shape_add(rank, 0);
    shape_add[pack_dim] = 1;
    shape_value = builder.create<TF::AddOp>(
        loc, shape_type, shape_value,
        CreateConstValue<int64_t>(builder, loc, {rank}, shape_add));
  }
  packed_shape[pack_dim] /= 2;
  SmallVector<int64_t> divisor(rank, 1);
  divisor[pack_dim] = 2;

  RankedTensorType packed_output_type =

// If <handle> has rank <rank>, or its rank is unknown, return OK and return the
// shape with asserted rank in <*result>. Otherwise an error is placed into
// `status`.
TF_CAPI_EXPORT extern void TF_ShapeInferenceContextWithRank(
    TF_ShapeInferenceContext* ctx, TF_ShapeHandle* handle, int64_t rank,
    TF_ShapeHandle* result, TF_Status* status);

// If <handle> has rank at least <rank>, or its rank is unknown, return OK and