add stablehlo pass to convert IR with quantized types to standard types (#24785)

On current main the test i added here fails with:

```
input_qdq.mlir:4:24: error: failed to legalize operation 'stablehlo.uniform_quantize' that was explicitly marked illegal: %0 = "stablehlo.uniform_quantize"(%arg0) : (tensor<1x256xf32>) -> tensor<1x256x!quant.uniform<i8:f32, 1.000000e-02>>
    %quantized_input = stablehlo.uniform_quantize %input : (tensor<1x256xf32>) -> tensor<1x256x!quant.uniform<i8:f32, 0.01:0>>
                       ^
input_qdq.mlir:4:24: note: see current operation: %0 = "stablehlo.uniform_quantize"(%arg0) : (tensor<1x256xf32>) -> tensor<1x256x!quant.uniform<i8:f32, 1.000000e-02>>
```

This PR adds a shlo pass to convert IR with quantized types to integer
math + standard types.

Signed-off-by: default <ziereis@roofline.ai>

compiler/plugins/input/StableHLO/Conversion/Passes.cpp

diff --git a/compiler/plugins/input/StableHLO/Conversion/Passes.cpp b/compiler/plugins/input/StableHLO/Conversion/Passes.cpp
index 3a0dd8e..ff5fc90 100644
--- a/compiler/plugins/input/StableHLO/Conversion/Passes.cpp
+++ b/compiler/plugins/input/StableHLO/Conversion/Passes.cpp
@@ -58,6 +58,10 @@
       ::mlir::stablehlo::createStablehloLegalizeCompositeToCallPass(
           ::mlir::stablehlo::StablehloLegalizeCompositeToCallPassOptions{
               /*exceptListOption=*/{}}));
+  passManager.addNestedPass<func::FuncOp>(
+      ::mlir::stablehlo::createStablehloLegalizeQuantToMathPass());
+  passManager.addNestedPass<func::FuncOp>(
+      ::mlir::stablehlo::createChloLegalizeToStablehloPass());
   passManager.addNestedPass<func::FuncOp>(mlir::createCanonicalizerPass());
   passManager.addNestedPass<func::FuncOp>(createStableHLOCanonicalize());
   passManager.addNestedPass<func::FuncOp>(mlir::createCSEPass());

compiler/plugins/input/StableHLO/Conversion/test/auto_input_conversion.mlir

diff --git a/compiler/plugins/input/StableHLO/Conversion/test/auto_input_conversion.mlir b/compiler/plugins/input/StableHLO/Conversion/test/auto_input_conversion.mlir
index efb9b89..3cd1fb3 100644
--- a/compiler/plugins/input/StableHLO/Conversion/test/auto_input_conversion.mlir
+++ b/compiler/plugins/input/StableHLO/Conversion/test/auto_input_conversion.mlir
@@ -11,6 +11,33 @@
 
 // -----
 
+// Check that quantized StableHLO is lowered to primitive integer/float math.
+
+// CHECK-LABEL: util.func public @qdq_linear
+// CHECK-SAME: %{{.*}}: tensor<1x256xf32>
+// CHECK-SAME: %{{.*}}: tensor<256x512xi8>
+// CHECK: linalg.vecmat
+// CHECK-NOT: !quant.uniform
+// CHECK-NOT: stablehlo.uniform_quantize
+// CHECK-NOT: stablehlo.uniform_dequantize
+func.func @qdq_linear(
+    %input: tensor<1x256xf32>,
+    %weight: tensor<256x512x!quant.uniform<i8:f32, 0.01:0>>
+  ) -> tensor<1x512xf32> {
+  %quantized_input = stablehlo.uniform_quantize %input :
+      (tensor<1x256xf32>) -> tensor<1x256x!quant.uniform<i8:f32, 0.01:0>>
+  %matmul_result = stablehlo.dot_general %quantized_input, %weight,
+      contracting_dims = [1] x [0] :
+      (tensor<1x256x!quant.uniform<i8:f32, 0.01:0>>,
+       tensor<256x512x!quant.uniform<i8:f32, 0.01:0>>) ->
+      tensor<1x512x!quant.uniform<i32:f32, 0.0001:0>>
+  %output = stablehlo.uniform_dequantize %matmul_result :
+      (tensor<1x512x!quant.uniform<i32:f32, 0.0001:0>>) -> tensor<1x512xf32>
+  return %output : tensor<1x512xf32>
+}
+
+// -----
+
 // CHECK-LABEL: util.func public @vhlo_func
 vhlo.func_v1 @vhlo_func(%arg0: !vhlo.tensor_v1<!vhlo.i32_v1>) -> (!vhlo.tensor_v1<!vhlo.i32_v1>) {
   // CHECK: arith.constant