// RUN: tf-opt -xla-legalize-tf %s | FileCheck %s

//===----------------------------------------------------------------------===//
// tf.BatchMatMulV2 op legalizations.
//===----------------------------------------------------------------------===//

func.func @batchmatmulv2_basic(%arg0: tensor<1x4x2xf32>, %arg1: tensor<3x2x4xf32>) -> tensor<3x4x4xf32> {
// CHECK-LABEL:   func @batchmatmulv2_basic

// CHECK-LABEL: batchmatmulv2
func.func @batchmatmulv2(%arg0: tensor<1x4x2xf32>, %arg1: tensor<3x2x4xf32>) -> tensor<3x4x4xf32> {
  // NO_FALLBACK: mhlo.dynamic_broadcast_in_dim
  // NO_FALLBACK: mhlo.dot_general

  // SUPPORTED_FALLBACK_DEVICE: mhlo.reduce
  // SUPPORTED_FALLBACK_DEVICE: mhlo.dot_general
  // SUPPORTED_FALLBACK_DEVICE: mhlo.transpose

  %0 = "tf.Einsum"(%arg0, %arg1) {T = "tfdtype$DT_FLOAT", equation = "ijk,ikm->ijm"}: (tensor<3x4x5xf32>, tensor<3x5x6xf32>) -> tensor<3x4x6xf32>
  func.return %0 : tensor<3x4x6xf32>
  // CHECK-LABEL: einsum_basic
  // CHECK: "tf.BatchMatMulV2"(%arg0, %arg1) <{adj_x = false, adj_y = false}> : (tensor<3x4x5xf32>, tensor<3x5x6xf32>) -> tensor<3x4x6xf32>
}

func.func @einsum_matmul(%arg0: tensor<7x9xf32>, %arg1: tensor<9x5xf32>) -> tensor<7x5xf32> {

  %0 = "tf.BroadcastTo"(%arg1, %cst) : (tensor<17x24xf32>, tensor<3xi64>) -> tensor<17x17x24xf32>
  %1 = "tf.BatchMatMulV2"(%arg0, %0) {adj_x = false, adj_y = false} : (tensor<17x17x17xf32>, tensor<17x17x24xf32>) -> tensor<17x17x24xf32>
  func.return %1 : tensor<17x17x24xf32>
  // CHECK: %[[V0:.*]] = "tf.BatchMatMulV2"(%arg0, %arg1) <{adj_x = false, adj_y = false}> : (tensor<17x17x17xf32>, tensor<17x24xf32>) -> tensor<17x17x24xf32>

  %0 = "tf.BatchMatMulV2"(%arg0, %arg3) {adj_x = false, adj_y = false} : (tensor<?x8x8xf32>, tensor<8x40xf32>) -> tensor<?x8x40xf32>
  %1 = "tf.Add"(%0, %arg5) : (tensor<?x8x40xf32>, tensor<40xf32>) -> tensor<?x8x40xf32>
  %2 = "tf.BatchMatMulV2"(%1, %arg4) {adj_x = false, adj_y = true} : (tensor<?x8x40xf32>, tensor<10x40xf32>) -> tensor<?x8x10xf32>

    key: "shape"
    value {
      shape {
        dim {
          size: 3
        }
        dim {
          size: 7
        }
      }
    }
  }
}
node {
  name: "MatMul"
  op: "BatchMatMulV2"
  input: "Placeholder"
  input: "Placeholder_1"
  attr {
    key: "T"
    value {
      type: DT_FLOAT
    }
  }
  attr {
    key: "adj_x"
    value {
      b: false
    }
  }

#include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h"
#include "tensorflow/core/util/matmul_bcast.h"

namespace mlir {
namespace TF {

// Populate patterns to unroll tf.BatchMatMulV2 op into a sequence of TF ops.
// Since TFLite does not support BatchMatMul operation, it unrolls a BatchMatMul
// op into tf.Reshape, tf.Slice, tf.MatMul, tf.Pack, and tf.Reshape ops.

    key: "shape"
    value {
      shape {
        dim {
          size: 3
        }
        dim {
          size: 7
        }
      }
    }
  }
}
node {
  name: "MatMul"
  op: "BatchMatMulV2"
  input: "Placeholder"
  input: "Placeholder_1"
  attr {
    key: "T"
    value {
      type: DT_FLOAT
    }
  }
  attr {
    key: "adj_x"
    value {
      b: false
    }
  }

  %cst = "tf.Const"() {device = "", value = dense<1.000000e+01> : tensor<10x2xbf16>} : () -> tensor<10x2xbf16>
  %0 = "tf.Cast"(%arg0) {Truncate = false, device = ""} : (tensor<1x1x10xf32>) -> tensor<1x1x10xbf16>
  %1 = "tf.BatchMatMulV2"(%0, %cst) {adj_x = false, adj_y = false, device = ""} : (tensor<1x1x10xbf16>, tensor<10x2xbf16>) -> tensor<1x1x2xbf16>
  %2 = "tf.Cast"(%1) {Truncate = false} : (tensor<1x1x2xbf16>) -> tensor<1x1x2xf32>

cluster 9 size 5
 All 1
 Const 1
 Less 1
 LogicalAnd 1
 LogicalNot 1
cluster 10 size 9
 All 1
 Const 4
 Equal 1
 LessEqual 1
 LogicalOr 1
 Max 1
cluster 11 size 302
 Add 24
 BatchMatMulV2 1
 BiasAdd 8
 Cast 8
 ConcatV2 16
 Const 81
 ExpandDims 3
 Fill 1
 GreaterEqual 8
 Identity 1
 Less 1
 MatMul 10
 Mul 44
 Range 1
 Rsqrt 1
 Select 19
 Shape 6