[Preprocessing] Handle batching dims and dynamic-update-slice scatter ops (#24720)

See bug report https://github.com/iree-org/iree/issues/24719

JAX-emitted scatters use two forms that IREE's StableHLO scatter
preprocessing did not support, causing malformed shape errors during
input conversion:

1. Scatter batching dims were only passed through, so the downstream
canonicalizers collapsed batching dims together with the real scatter
loop dim and miscompiled. Added `ScatterBatchingDimsExpander` (benefit
2). This is ported with just a few style tweaks directly from upstream
StableHLO's `ScatterWithBatchingDimsExpander`. I verified that the
StableHLO pattern wasn't accessible to be called directly. The only way
to access it is to run the whole set of compatibility expanders which
seemed brittle since new expanders can be added over time, so just
copied it and its helper functions. I didn't spend time to fully
understand the code though since it's an already reviewed pattern used
to handle this exact case for backwards compatibility in StableHLO.
Basically, it lowers each batching dim to an explicit iota index column
so no batching dims remain for the existing patterns. Here's an example
```mlir
module @m {
  func.func @s(%operand: tensor<8x33xi32>, %indices: tensor<8x32x1xi32>, %updates: tensor<8x32xi32>) -> tensor<8x33xi32> {
    %0 = "stablehlo.scatter"(%operand, %indices, %updates) <{scatter_dimension_numbers = #stablehlo.scatter<inserted_window_dims = [1], input_batching_dims = [0], scatter_indices_batching_dims = [0], scatter_dims_to_operand_dims = [1], index_vector_dim = 2>}> ({
    ^bb0(%arg2: tensor<i32>, %arg3: tensor<i32>):
      %1 = stablehlo.minimum %arg2, %arg3 : tensor<i32>
      stablehlo.return %1 : tensor<i32>
    }) : (tensor<8x33xi32>, tensor<8x32x1xi32>, tensor<8x32xi32>) -> tensor<8x33xi32>
    return %0 : tensor<8x33xi32>
  }
}
```
gets preprocessed into
```mlir
module @m {
  func.func @s(%arg0: tensor<8x33xi32>, %arg1: tensor<8x32x1xi32>, %arg2: tensor<8x32xi32>) -> tensor<8x33xi32> {
    %0 = stablehlo.iota dim = 0 : tensor<8x32x1xi32>
    %1 = stablehlo.concatenate %0, %arg1, dim = 2 : (tensor<8x32x1xi32>, tensor<8x32x1xi32>) -> tensor<8x32x2xi32>
    %collapsed = tensor.collapse_shape %1 [[0, 1], [2]] : tensor<8x32x2xi32> into tensor<256x2xi32>
    %collapsed_0 = tensor.collapse_shape %arg2 [[0, 1]] : tensor<8x32xi32> into tensor<256xi32>
    %2 = "stablehlo.scatter"(%arg0, %collapsed, %collapsed_0) <{indices_are_sorted = false, scatter_dimension_numbers = #stablehlo.scatter<inserted_window_dims = [0, 1], scatter_dims_to_operand_dims = [0, 1], index_vector_dim = 1>, unique_indices = false}> ({
    ^bb0(%arg3: tensor<i32>, %arg4: tensor<i32>):
      %3 = stablehlo.minimum %arg3, %arg4 : tensor<i32>
      stablehlo.return %3 : tensor<i32>
    }) : (tensor<8x33xi32>, tensor<256x2xi32>, tensor<256xi32>) -> tensor<8x33xi32>
    return %2 : tensor<8x33xi32>
  }
}
```

2. Single-index, full-rank overwrite scatters are dynamic-update-slice
semantics and cannot be represented as `iree_linalg_ext.scatter`. Added
`ScatterToDynamicUpdateSlice` (benefit 3) to rewrite them into
`stablehlo.dynamic_update_slice`.
Here's an example:
```mlir
module @m {
  func.func @s(%operand: tensor<8x7x24x32xf32>, %indices: tensor<1xi32>, %updates: tensor<8x4x24x32xf32>) -> tensor<8x7x24x32xf32> {
    %0 = "stablehlo.scatter"(%operand, %indices, %updates) <{indices_are_sorted = true, scatter_dimension_numbers = #stablehlo.scatter<update_window_dims = [0, 1, 2, 3], scatter_dims_to_operand_dims = [1]>, unique_indices = true}> ({
    ^bb0(%arg2: tensor<f32>, %arg3: tensor<f32>):
      stablehlo.return %arg3 : tensor<f32>
    }) : (tensor<8x7x24x32xf32>, tensor<1xi32>, tensor<8x4x24x32xf32>) -> tensor<8x7x24x32xf32>
    return %0 : tensor<8x7x24x32xf32>
  }
}
```
gets preprocessed into
```mlir
module @m {
  func.func @s(%arg0: tensor<8x7x24x32xf32>, %arg1: tensor<1xi32>, %arg2: tensor<8x4x24x32xf32>) -> tensor<8x7x24x32xf32> {
    %c = stablehlo.constant dense<0> : tensor<i32>
    %0 = stablehlo.slice %arg1 [0:1] : (tensor<1xi32>) -> tensor<1xi32>
    %1 = stablehlo.reshape %0 : (tensor<1xi32>) -> tensor<i32>
    %2 = stablehlo.dynamic_update_slice %arg0, %arg2, %c, %1, %c, %c : (tensor<8x7x24x32xf32>, tensor<8x4x24x32xf32>, tensor<i32>, tensor<i32>, tensor<i32>, tensor<i32>) -> tensor<8x7x24x32xf32>
    return %2 : tensor<8x7x24x32xf32>
  }
}
```

Both patterns run before the existing scatter canonicalizers, because
the first expands out functionality not supported by later patterns, and
the second one removes the scatter entirely. Lit tests added for both.

---------

Signed-off-by: Paul Stark <paul.stark@cdprojektred.com>
Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>