REGISTER_KERNEL_BUILDER(Name("XlaVariadicReduceV2").Device(DEVICE),          \
                          XlaCompileOnDemandOp);                               \
  REGISTER_KERNEL_BUILDER(Name("XlaReducePrecision").Device(DEVICE),           \
                          XlaCompileOnDemandOp);                               \
  REGISTER_KERNEL_BUILDER(Name("XlaReduceWindow")                              \

      "XlaEinsum",
      "XlaGather",
      "XlaIf",
      "XlaKeyValueSort",
      "XlaOptimizationBarrier",
      "XlaPad",
      "XlaRecv",
      "XlaReduce",
      "XlaReducePrecision",
      "XlaReduceScatter",
      "XlaReduceWindow",
      "XlaRemoveDynamicDimensionSize",
      "XlaReplicaId",
      "XlaRngBitGenerator",
      "XlaScatter",
      "XlaSelectAndScatter",

    BuildBodyWithCall(rewriter, loc, func, func_ty, &sort_op.getComparator());

    rewriter.replaceOp(op, sort_op.getResults());
    return success();
  }
};

// Convert tf.XlaReducePrecision to mhlo.ReducePrecision
class ConvertXlaReducePrecisionOp
    : public OpRewritePattern<TF::XlaReducePrecisionOp> {
 public:
  using OpRewritePattern::OpRewritePattern;

  );

  TF_DerivedOperandTypeAttr T = TF_DerivedOperandTypeAttr<0>;

  let hasCanonicalizer = 1;
}

def TF_XlaReducePrecisionOp : TF_Op<"XlaReducePrecision", [Pure]> {
  let summary = "Wraps the XLA ReducePrecision operator";

  let description = [{
documented at https://www.tensorflow.org/xla/operation_semantics#reduceprecision.
  }];