// RUN: tf-opt -split-input-file -verify-diagnostics -xla-legalize-tf-collective -xla-legalize-tf %s | FileCheck %s


// CHECK: module attributes
// CHECK-SAME{LITERAL}: tf2xla.collective_info.group_key = 0
// CHECK-SAME{LITERAL}: tf2xla.collective_info.group_size = 2
// CHECK-LABEL: func @all_reduce_cross_replica
func.func @all_reduce_cross_replica(%input: tensor<f32>) -> tensor<f32> {

// (unlike the rest), since this is the primary use case for such ops and
// verification of shapes and broadcasts is desired.
// RUN: tf-opt "-xla-legalize-tf=legalize-chlo=true" -canonicalize %s | FileCheck %s
// RUN: tf-opt "-xla-legalize-tf=legalize-chlo=false" %s | FileCheck --check-prefix CHLO %s

//===----------------------------------------------------------------------===//
// Binary op legalizations.

def LegalizeRelu6 : Pat<(TF_Relu6Op $arg), (TFL_Relu6Op $arg)>;
def LegalizeRelu : Pat<(TF_ReluOp $arg), (TFL_ReluOp $arg)>;
// TFL Relu doesn't support I32/I64 type, so legalizes TF Relu to TFL Maximum.
def LegalizeReluI32 :
  Pat<(TF_ReluOp TensorOf<[I32]>:$arg),
      (TFL_MaximumOp $arg,
        (Arith_ConstantOp ConstantAttr<RankedI32ElementsAttr<[]>,"0">))>;
def LegalizeReluI64 :

// RUN: tf-opt -convert-tf-quant-types -quant-convert-tf-quant-ops-to-mhlo -canonicalize "-xla-legalize-tf=legalize-chlo=false" -split-input-file %s | FILECHECK_OPTS="" FileCheck %s


//===----------------------------------------------------------------------===//
// tf.UniformQuantizedDotHybrid legalization
//===----------------------------------------------------------------------===//

// CHECK-LABEL: func @quantized_matmul_fn

// RUN: tf-opt -split-input-file -verify-diagnostics -xla-legalize-tf-communication %s | FileCheck %s

// Test legalization of `tf._XlaHostComputeMlir` expands into individual
// `mhlo.send` per operand and `mhlo.recv` per result. Channel Id's are uniquely
// assigned per mhlo communcation op, and frontend attributes (modified keys)
// and op shardings are added. Sink tokens are created
// if there are more than one operand or more than one result.
//

    }

    def "cannot add elements after collection is realized"() {
        given:
        container.add(Stub(SwiftBinary))
        container.realizeNow()

        when:
        container.add(Stub(SwiftBinary))

        then:
        def e = thrown(IllegalStateException)
        e.message == 'Cannot add an element to this collection as it has already been realized.'
    }

// RUN: tf-opt "-xla-legalize-tf=device-type=XLA_CPU_JIT prefer-tf2xla=true use-tf2xla-fallback=true" %s -verify-diagnostics -mlir-disable-threading  | FileCheck %s

module attributes {tf.versions = {bad_consumers = [], min_consumer = 0 : i32, producer = 268 : i32}} {
  // CHECK-LABEL: binary_op
  func.func @binary_op(%arg0: tensor<2xf32>, %arg1: tensor<2xf32>) -> tensor<2xf32> {
    // CHECK: mhlo.atan2

    pass_manager->addPass(mlir::TF::CreateTFShapeInferencePass());
  }

  // Legalize while early to allow further constant folding.
  // TODO(jpienaar): This may not actually matter as we do canonicalization
  // after the legalize below, for now it needs to be below the above passes
  // that work on TF dialect and before inliner so that the function calls in

        }

        @Override
        public void collectInto(ImmutableCollection.Builder<T> builder) {
            if (!realized) {
                realize();
            }
            builder.addAll(cache);
        }

        List<T> getValues() {
            if (!realized) {
                realize();
            }
            return cache;
        }

        public boolean remove(T value) {

          for (auto arg : block.getArguments()) {
            ConvertAndWrapUsesInUnrealizedCast(arg, converter, rewriter);
          }
        }
      }
    });
  });
  return success();
}

// Legalize StableHLO portion of program to VHLO, leaves TFL untouched
LogicalResult ApplyStablehloToVhloPatterns(ModuleOp module,
                                           bool is_func_legal) {
  MLIRContext *context = module.getContext();