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

                                         "XlaSpmdShardToFullShape",
                                         "XlaSvd",
                                         "XlaVariadicReduceV2",
                                         "XlaVariadicSort",
                                         "XlaWhile"};

static bool NodeCanTriggerXlaCompilation(const NodeDef& node) {

}


// CHECK-LABEL: testXlaReduceToXlaVariadicReduceV2
func.func @testXlaReduceToXlaVariadicReduceV2(%arg0: tensor<*xbf16>, %arg1: tensor<*xbf16>) -> tensor<*xbf16> {
  // CHECK: "tf.XlaVariadicReduceV2"(%arg0, %arg1) <{dimensions_to_reduce = [], operandSegmentSizes = array<i32: 1, 1>, reducer = @sum1}> {device = "/job:localhost/replica:0/task:0/device:GPU:0"} : (tensor<*xbf16>, tensor<*xbf16>) -> tensor<*xbf16>

      "XlaSharding",
      "XlaSort",
      "XlaSplitND",
      "XlaSpmdFullToShardShape",
      "XlaSpmdShardToFullShape",
      "XlaSvd",
      "XlaVariadicReduce",
      "XlaVariadicReduceV2",
      "XlaVariadicSort",
      "XlaWhile",
      "Zeta",
      "_Arg",
      "_ArrayToList",
      "_ListToArray",
      "_Retval"};
  return result;
}

      }
    }

    if (op.getDimensionsToReduce().size() > input_ty_0.getRank()) {
      return op.emitOpError()
             << "Invalid dimensions_to_reduce argument to XlaVariadicReduceV2";
    }
  }

  for (int i = 0; i < n_inputs; ++i) {
    auto init_value_ty_i = init_values_ty[i].cast<ShapedType>();
    if (init_value_ty_i.hasRank() && init_value_ty_i.getRank() != 0) {

  return ::xla::ConvertPrecisionConfig(&precision, builder);
}

//===----------------------------------------------------------------------===//
// XlaVariadicReduceV2 op utilities.
//===----------------------------------------------------------------------===//

static void BuildBodyWithCall(PatternRewriter &rewriter, const Location &loc,

  // CHECK-NEXT:   %[[SUM:.*]] = func.call @sum_reducer(%[[ARG0]], %[[ARG1]]){{.*}}
  // CHECK-NEXT:   mhlo.return %[[SUM]] : tensor<complex<f64>>
  // CHECK: return %[[REDUCE]]

  let description = [{
Semantics are documented at
 https://www.tensorflow.org/performance/xla/operation_semantics#variadic_reduce.

This version is limited to operands of the same dtype.
XlaVariadicReduceV2 is a version that supports heterogeneous operands.
  }];

  let arguments = (ins