// CHECK-NOT: BroadcastTo
    // CHECK: TensorScatterUpdate
    "tf.ResourceScatterUpdate"(%resource, %indices, %updates) {device = ""} : (tensor<*x!tf_type.resource>, tensor<?xi32>, tensor<i32>) -> ()
    tf_device.return
  }) : () -> ()
  func.return
}

// Test that composite tf.ResourceScatterUpdate operation is decomposed to tf.TensorScatterUpdate when update is unranked
// CHECK-LABEL: @decompose_unranked_updates

       (TF_ConstOp (GetScalarOfType<-1> $indices))),
      $updates
    ),
    (CreateConstBoolAttrFalse)
  )>;

// Pattern to decompose tf.ResourceScatterUpdate into tf.ReadVariable,
// tf.TensorScatterUpdate, and tf.AssignVariable.
def DecomposeResourceScatterUpdate : Pat<
  (TF_ResourceScatterUpdateOp:$src_op $resource, $indices, $updates),
  (TF_AssignVariableOp
    $resource,
    (TF_TensorScatterUpdateOp

                          /*incompatible_shape_error*/ConstBoolAttrTrue)>;

//===----------------------------------------------------------------------===//
// TensorScatterUpdate op patterns.
//===----------------------------------------------------------------------===//

def LowerTensorScatterUpdate_1 : Pat<(TF_TensorScatterUpdateOp $input, $indices, $updates),
          (TF_CastOp

  if (maybe_contiguous_start == 0 && buffer_type == updates_type) {
    return AccumulateBuffers(buffer, updates, builder, loc);
  }
  // We cannot simply use a TensorScatterUpdate, as it does not accumulate with
  // the old data; it is tricky to manually add the old data either, since there
  // could be duplicates in the index. We follow the old bridge's approach by

                                      nudged_float_min);

    rewriter.replaceOp(op, {output});
    return success();
  }
};

// Lowers InvertPermutation op to TensorScatterUpdate op.
//
// Example:
//
//   %x = "tf.Const"() {value = dense<[3, 4, 0, 1, 2]> : tensor<5xi32>}
//   "tf.InvertPermutation"(%x) : (tensor<5xi32>) -> tensor<5xi32>
//
// is lowered to
//

  // CHECK:  %[[ADD:.*]] = "tf.Add"(%arg2, %[[GATHER]]) : (tensor<1xi64>, tensor<1xi64>) -> tensor<1xi64>
  // CHECK:  %[[SCATTER:.*]] = "tf.TensorScatterUpdate"(%arg0, %arg1, %[[ADD]]) : (tensor<7xi64>, tensor<1x1xi32>, tensor<1xi64>) -> tensor<7xi64>
  // CHECK:  return %[[SCATTER]] : tensor<7xi64>
}

  // CHECK-SAME: unique_indices = false
  // CHECK:  ^bb0(%arg3: tensor<f32>, %arg4: tensor<f32>):
  // CHECK:    mhlo.return %arg4 : tensor<f32>
  // CHECK:  })
  %0 = "tf.TensorScatterUpdate"(%tensor, %indices, %updates) : (tensor<?x?x?xf32>, tensor<?x2xi32>, tensor<?x?xf32>) -> tensor<?x?x?xf32>
  func.return %0 : tensor<?x?x?xf32>
}

// -----

                                     hlo::kInfinityMax, &rewriter);
  }

  static ComparisonDirection GetDirection() { return ComparisonDirection::LE; }
};

// Converts TF TensorScatterUpdate/Min/Max/Add/Sub op into Scatter Op with
// assignment:
//
//   %result = "mhlo.scatter"(%tensor, %indices, %updates)
//     { dimensions = ... }
//
template <typename Derived, typename OpTy>

  );

  TF_DerivedOperandTypeAttr T = TF_DerivedOperandTypeAttr<0>;
  TF_DerivedOperandTypeAttr Tindices = TF_DerivedOperandTypeAttr<1>;
}

def TF_TensorScatterUpdateOp : TF_Op<"TensorScatterUpdate", [Pure]> {
  let summary = [{
Scatter `updates` into an existing tensor according to `indices`.
  }];

  let description = [{