From d5042e75890feefbccfa9cce406d8a6344ca25f8 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Fri, 23 Feb 2024 19:05:42 +0000 Subject: [PATCH 01/14] Quantization Verifiers : T2x constraits --- stablehlo/dialect/Base.cpp | 39 +- stablehlo/dialect/StablehloOps.td | 4 +- stablehlo/tests/ops_stablehlo.mlir | 10 +- stablehlo/tests/ops_stablehlo_quantized.mlir | 2 +- .../stablehlo_legalize_to_vhlo.0_10_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_11_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_12_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_13_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_14_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_15_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_16_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_17_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_18_0.mlir | 8 +- .../stablehlo_legalize_to_vhlo.0_9_0.mlir | 8 +- .../vhlo/stablehlo_legalize_to_vhlo.mlir | 8 +- t2.patch | 1105 +++++++++++++++++ 16 files changed, 1175 insertions(+), 73 deletions(-) create mode 100644 t2.patch diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index 9d0fcb5c274..ec0babf36eb 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -72,31 +72,34 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { tp1 = getElementTypeOrSelf(tp1); tp2 = getElementTypeOrSelf(tp2); - // Quantization: In the most general case, we allow any combination of - // quantized/non-quantized across any combination of operands/results, - // and some differences in quantization parameters across operands/results. - // Individual ops may introduce additional constraints. auto qtp1 = tp1.dyn_cast(); auto qtp2 = tp2.dyn_cast(); + + if (!qtp1 && !qtp2) { + // Both are non quantized + return tp1 == tp2; + } if (qtp1 && qtp2) { + // Both are quantized if (qtp1.getStorageType() != qtp2.getStorageType() || qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || - qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) + qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax() || + qtp1.getExpressedType() != qtp2.getExpressedType()){ return false; + } + + auto qpatp1 = qtp1.dyn_cast(); + auto qpatp2 = qtp2.dyn_cast(); + if(qpatp1 && qpatp2){ + // Both are also per-axis quantized + // For now, don't match dimentions, Per OP verifier will do it. + return true; + } + // return true if both are per-tensor quantized + return !(qpatp1 || qpatp2); } - auto etp1 = getExpressedTypeOrSelf(tp1); - auto etp2 = getExpressedTypeOrSelf(tp2); - - // Sparsity: In the most general case, we allow any combination of - // sparsity/denseness across any combination of operands/results, as well as - // differences in sparsity encodings for operands and results. - // Individual ops may introduce additional constraints. - // No additional code is needed to check this because of how sparsity is - // currently implemented. - - // Default case: Unless dynamism, quantization and/or sparsity are involved, - // the types are required to be exactly equal. - return etp1 == etp2; + + return false; } bool isCompatibleForHloTypeInference(Type tp1, Type tp2) { diff --git a/stablehlo/dialect/StablehloOps.td b/stablehlo/dialect/StablehloOps.td index c8203897343..a1cf4f35466 100644 --- a/stablehlo/dialect/StablehloOps.td +++ b/stablehlo/dialect/StablehloOps.td @@ -1927,7 +1927,7 @@ def StableHLO_DynamicBroadcastInDimOp : StableHLO_ShapedInterfaceOp< // directly. def StableHLO_CholeskyOp : StableHLO_Op<"cholesky", - [Pure, SameOperandsAndResultElementType /*cholesky_c1*/, + [Pure, HLO_CompatibleOperandsAndResultElementType /*cholesky_c1*/, InferTensorType /*cholesky_c1*/]> { let summary = "Cholesky operation"; let description = [{ @@ -1954,7 +1954,7 @@ def StableHLO_CholeskyOp : StableHLO_Op<"cholesky", } def StableHLO_ClampOp : StableHLO_ShapedInterfaceOp<"clamp", [Pure, - SameOperandsAndResultElementType /* clamp_c3 */, HLO_BroadcastingElementwise, + HLO_CompatibleOperandsAndResultElementType /* clamp_c3 */, HLO_BroadcastingElementwise, InferTensorType]> { let summary = "Clamp operation"; let description = [{ diff --git a/stablehlo/tests/ops_stablehlo.mlir b/stablehlo/tests/ops_stablehlo.mlir index f0d857cc8bb..96e65d84bc3 100644 --- a/stablehlo/tests/ops_stablehlo.mlir +++ b/stablehlo/tests/ops_stablehlo.mlir @@ -5070,16 +5070,10 @@ func.func @is_compatible_dynamism_dim_mismatch(%arg0: tensor<1x?xf32>) { // ----- -// TODO(b/230263270): For stablehlo.add, the plan is to only allow fp+fp=fp, q+q=q and q+q=fp. func.func @is_compatible_quant_mix_non_quant(%arg0: tensor<1xf32>, %arg1: tensor<1x!quant.uniform>) { %0 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32> - %1 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1x!quant.uniform> - %2 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> - %3 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> - %4 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> - %5 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> - %6 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> - %7 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> + %1 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> + %2 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> func.return } diff --git a/stablehlo/tests/ops_stablehlo_quantized.mlir b/stablehlo/tests/ops_stablehlo_quantized.mlir index f3dae90b165..d62a0d2e0bc 100644 --- a/stablehlo/tests/ops_stablehlo_quantized.mlir +++ b/stablehlo/tests/ops_stablehlo_quantized.mlir @@ -84,7 +84,7 @@ func.func @ops_per_tensor_quantization( %sqrt = "stablehlo.sqrt"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> %subtract = "stablehlo.subtract"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> %tanh = "stablehlo.tanh"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> - %transpose = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> + %transpose = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform:f32, 1.0:17>> %tuple = "stablehlo.tuple"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tuple>, tensor<1x2x2x!quant.uniform>> %uniform_dequantize = "stablehlo.uniform_dequantize" (%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2xf32> %uniform_quantize = "stablehlo.uniform_quantize" (%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir index 57bb204fa3c..8fa17f8efb3 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir @@ -2226,10 +2226,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir index 46f5321d57b..f97ca3e6d8b 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir @@ -2233,10 +2233,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir index 816d570de06..5889c4a4cdc 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir @@ -2233,10 +2233,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir index 9616dac9456..4ff99bb1612 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir @@ -2233,10 +2233,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir index 49195e7ab94..981b356b868 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir @@ -2233,10 +2233,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir index 1902ad40f83..578dffe926a 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir @@ -2241,10 +2241,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir index a2bbf19b2f2..17a062b567b 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir @@ -2252,10 +2252,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir index c8594b43b89..77acf931740 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir @@ -2376,10 +2376,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir index 65a7412b53e..8aa808cc1e6 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir @@ -2376,10 +2376,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_per_tensor_quantization" -func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK-LABEL: "type_per_axis_quantization" diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir index 4036e171206..fbd4121b57b 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir @@ -2212,10 +2212,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_quantization" -func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir index 36bf834f882..5559f6ca662 100644 --- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir +++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir @@ -2381,10 +2381,10 @@ func.func @type_dynamism_ranked(%arg0: tensor) -> tensor { } // CHECK-LABEL: "type_per_tensor_quantization" -func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { - // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 - %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor - func.return %0 : tensor +func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { + // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> + %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> + func.return %0 : tensor> } // CHECK-LABEL: "type_per_axis_quantization" diff --git a/t2.patch b/t2.patch new file mode 100644 index 00000000000..5437a25b1ff --- /dev/null +++ b/t2.patch @@ -0,0 +1,1105 @@ +From cb8d9384082235518fb6c750b73959503400c740 Mon Sep 17 00:00:00 2001 +From: Abhinav +Date: Wed, 21 Feb 2024 19:46:10 +0000 +Subject: [PATCH 1/5] T2 Trait verifier + +--- + stablehlo/dialect/Base.cpp | 11 ++++++++++- + stablehlo/dialect/StablehloOps.td | 10 +++++----- + stablehlo/tests/ops_stablehlo.mlir | 4 ++-- + 3 files changed, 17 insertions(+), 8 deletions(-) + +diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp +index 9eb4f2340e..88b421b2d9 100644 +--- a/stablehlo/dialect/Base.cpp ++++ b/stablehlo/dialect/Base.cpp +@@ -83,7 +83,16 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || + qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) + return false; +- } ++ ++ auto qpatp1 = tp1.dyn_cast(); ++ auto qpatp2 = tp2.dyn_cast(); ++ if (qpatp1 && qpatp2) { ++ if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) ++ return false; ++ } ++ } ++ //if((qtp1 == NULL && qtp2 != NULL) || (qtp1 != NULL && qtp2 == NULL)) return false; ++ + auto etp1 = getExpressedTypeOrSelf(tp1); + auto etp2 = getExpressedTypeOrSelf(tp2); + +diff --git a/stablehlo/dialect/StablehloOps.td b/stablehlo/dialect/StablehloOps.td +index b2ab6d6389..29cdac7f95 100644 +--- a/stablehlo/dialect/StablehloOps.td ++++ b/stablehlo/dialect/StablehloOps.td +@@ -1847,7 +1847,7 @@ def StableHLO_BroadcastOp : StableHLO_ShapedInterfaceOp<"broadcast", + } + + def StableHLO_BroadcastInDimOp : StableHLO_Op<"broadcast_in_dim", +- [Pure, HLO_CompatibleOperandsAndResultElementType /*broadcast_in_dim_c1*/]> { ++ [Pure, /*can't use the Trait, need updates to verifierHLO_CompatibleOperandsAndResultElementType*/ /*broadcast_in_dim_c1*/]> { + let summary = "BroadcastInDim operation"; + let description = [{ + Expands the dimensions and/or rank of an input tensor by duplicating the +@@ -1926,7 +1926,7 @@ def StableHLO_DynamicBroadcastInDimOp : StableHLO_ShapedInterfaceOp< + // directly. + + def StableHLO_CholeskyOp : StableHLO_Op<"cholesky", +- [Pure, SameOperandsAndResultElementType /*cholesky_c1*/, ++ [Pure, HLO_CompatibleOperandsAndResultElementType /*cholesky_c1*/, + InferTensorType /*cholesky_c1*/]> { + let summary = "Cholesky operation"; + let description = [{ +@@ -1953,7 +1953,7 @@ def StableHLO_CholeskyOp : StableHLO_Op<"cholesky", + } + + def StableHLO_ClampOp : StableHLO_ShapedInterfaceOp<"clamp", [Pure, +- SameOperandsAndResultElementType /* clamp_c3 */, HLO_BroadcastingElementwise, ++ HLO_CompatibleOperandsAndResultElementType /* clamp_c3 */, HLO_BroadcastingElementwise, + InferTensorType]> { + let summary = "Clamp operation"; + let description = [{ +@@ -2455,7 +2455,7 @@ def StableHLO_MapOp: StableHLO_ShapedInterfaceOp<"map", + } + + def StableHLO_ReshapeOp: StableHLO_Op<"reshape", +- [Pure, HLO_CompatibleOperandsAndResultElementType]> { ++ [Pure, /* can't use the Trait, need verifier updatesHLO_CompatibleOperandsAndResultElementType*/]> { + let summary = "Reshape operation"; + let description = [{ + Performs reshape of `operand` tensor to a `result` tensor. +@@ -2784,7 +2784,7 @@ def StableHLO_TraceOp: StableHLO_Op<"trace", []> { + } + + def StableHLO_TransposeOp: StableHLO_ShapedInterfaceOp<"transpose", +- [Pure, HLO_CompatibleOperandsAndResultElementType, ++ [Pure, /* can't use the Trait, need verifier to handle per-axis HLO_CompatibleOperandsAndResultElementType,*/ + DeclareOpInterfaceMethods]> { + let summary = "Transpose operation"; + let description = [{ +diff --git a/stablehlo/tests/ops_stablehlo.mlir b/stablehlo/tests/ops_stablehlo.mlir +index 998054478f..d5cabe763e 100644 +--- a/stablehlo/tests/ops_stablehlo.mlir ++++ b/stablehlo/tests/ops_stablehlo.mlir +@@ -5476,8 +5476,8 @@ func.func @quantization_supported_ops(%arg0: tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>, %arg1: tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) { + %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> + %1 = "stablehlo.broadcast_in_dim"(%arg1) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) -> tensor<2x2x2x!quant.uniform:f32:0, {0.1:-30, 0.1:-30}>> +- %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> +- %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> ++ // %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> ++ // %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> + func.return + } + + +From cc743f716ff9d41db7a6dca6ede0ab20e483c155 Mon Sep 17 00:00:00 2001 +From: Abhinav +Date: Wed, 21 Feb 2024 20:34:04 +0000 +Subject: [PATCH 2/5] reorg isCompatibleElementTypeForHloTypeInference + +--- + stablehlo/dialect/Base.cpp | 62 ++++++++++++++++++++------------------ + 1 file changed, 33 insertions(+), 29 deletions(-) + +diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp +index 88b421b2d9..0ebcc01cba 100644 +--- a/stablehlo/dialect/Base.cpp ++++ b/stablehlo/dialect/Base.cpp +@@ -66,7 +66,9 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { + } + return success(); + } +- ++// tp1 tp2 Result ++// Quantized Non Quantized false ++// .... + bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + // Get element type if shaped + tp1 = getElementTypeOrSelf(tp1); +@@ -78,34 +80,36 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + // Individual ops may introduce additional constraints. + auto qtp1 = tp1.dyn_cast(); + auto qtp2 = tp2.dyn_cast(); +- if (qtp1 && qtp2) { +- if (qtp1.getStorageType() != qtp2.getStorageType() || +- qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || +- qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) +- return false; +- +- auto qpatp1 = tp1.dyn_cast(); +- auto qpatp2 = tp2.dyn_cast(); +- if (qpatp1 && qpatp2) { +- if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) +- return false; +- } +- } +- //if((qtp1 == NULL && qtp2 != NULL) || (qtp1 != NULL && qtp2 == NULL)) return false; +- +- auto etp1 = getExpressedTypeOrSelf(tp1); +- auto etp2 = getExpressedTypeOrSelf(tp2); +- +- // Sparsity: In the most general case, we allow any combination of +- // sparsity/denseness across any combination of operands/results, as well as +- // differences in sparsity encodings for operands and results. +- // Individual ops may introduce additional constraints. +- // No additional code is needed to check this because of how sparsity is +- // currently implemented. +- +- // Default case: Unless dynamism, quantization and/or sparsity are involved, +- // the types are required to be exactly equal. +- return etp1 == etp2; ++ if(!(qtp1 && qtp2)){ ++ // one Q and another non Q ++ if (qtp1 || qtp2) return false; ++ ++ // Sparsity: In the most general case, we allow any combination of ++ // sparsity/denseness across any combination of operands/results, as well as ++ // differences in sparsity encodings for operands and results. ++ // Individual ops may introduce additional constraints. ++ // No additional code is needed to check this because of how sparsity is ++ // currently implemented. ++ ++ // Default case: Unless dynamism, quantization and/or sparsity are involved, ++ // the types are required to be exactly equal. ++ ++ auto etp1 = getExpressedTypeOrSelf(tp1); ++ auto etp2 = getExpressedTypeOrSelf(tp2); ++ return etp1 == etp2; ++ } ++ if (qtp1.getStorageType() != qtp2.getStorageType() || ++ qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || ++ qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) ++ return false; ++ ++ auto qpatp1 = tp1.dyn_cast(); ++ auto qpatp2 = tp2.dyn_cast(); ++ if (qpatp1 && qpatp2) { ++ if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) ++ return false; ++ } ++ return true; + } + + bool isCompatibleForHloTypeInference(Type tp1, Type tp2) { + +From ac17026d5de29f30d94931645c3efe7d28cd506c Mon Sep 17 00:00:00 2001 +From: Abhinav +Date: Thu, 22 Feb 2024 18:58:13 +0000 +Subject: [PATCH 3/5] updated isCompatibleElementTypeForHloTypeInference logic + +--- + stablehlo/dialect/Base.cpp | 67 ++++++++++++++++++-------------------- + 1 file changed, 32 insertions(+), 35 deletions(-) + +diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp +index 0ebcc01cba..998dd183eb 100644 +--- a/stablehlo/dialect/Base.cpp ++++ b/stablehlo/dialect/Base.cpp +@@ -66,50 +66,47 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { + } + return success(); + } +-// tp1 tp2 Result +-// Quantized Non Quantized false +-// .... ++ ++// Q: Quantized (per-tensor or per-axis) ++// NQ: Non Quantized ++// tp1 tp2 Result ++// NQ NQ true/false ++// NQ Q false ++// Q(per-tensor) Q(per-tensor) true/false ++// Q(per-tensor) Q(per-axis) false ++// Q(per-axis) Q(per-axis) true/false ++ + bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + // Get element type if shaped + tp1 = getElementTypeOrSelf(tp1); + tp2 = getElementTypeOrSelf(tp2); + +- // Quantization: In the most general case, we allow any combination of +- // quantized/non-quantized across any combination of operands/results, +- // and some differences in quantization parameters across operands/results. +- // Individual ops may introduce additional constraints. + auto qtp1 = tp1.dyn_cast(); + auto qtp2 = tp2.dyn_cast(); +- if(!(qtp1 && qtp2)){ +- // one Q and another non Q +- if (qtp1 || qtp2) return false; +- +- // Sparsity: In the most general case, we allow any combination of +- // sparsity/denseness across any combination of operands/results, as well as +- // differences in sparsity encodings for operands and results. +- // Individual ops may introduce additional constraints. +- // No additional code is needed to check this because of how sparsity is +- // currently implemented. +- +- // Default case: Unless dynamism, quantization and/or sparsity are involved, +- // the types are required to be exactly equal. +- +- auto etp1 = getExpressedTypeOrSelf(tp1); +- auto etp2 = getExpressedTypeOrSelf(tp2); +- return etp1 == etp2; ++ ++ // If both are non quantized ++ if (!qtp1 && !qtp2) { ++ return tp1 == tp2; + } +- if (qtp1.getStorageType() != qtp2.getStorageType() || +- qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || +- qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) +- return false; +- +- auto qpatp1 = tp1.dyn_cast(); +- auto qpatp2 = tp2.dyn_cast(); +- if (qpatp1 && qpatp2) { +- if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) +- return false; ++ // If both are per-tensor Quantized ++ if (qtp1 && qtp2) { ++ if (qtp1.getStorageType() != qtp2.getStorageType() || ++ qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || ++ qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax() || ++ qtp1.getExpressedType() != qtp2.getExpressedType()){ ++ return false; ++ } ++ ++ auto qpatp1 = qtp1.dyn_cast(); ++ auto qpatp2 = qtp2.dyn_cast(); ++ // If both are also per-axis quantized ++ if(qpatp1 && qpatp2){ ++ return qpatp1.getQuantizedDimension() == qpatp2.getQuantizedDimension(); ++ } ++ return !(qpatp1 || qpatp2); + } +- return true; ++ ++ return false; + } + + bool isCompatibleForHloTypeInference(Type tp1, Type tp2) { + +From 7c342ce16fda44dc20313b05dd47120a0b860e39 Mon Sep 17 00:00:00 2001 +From: Abhinav +Date: Fri, 23 Feb 2024 00:38:54 +0000 +Subject: [PATCH 4/5] generated new vhlo bytecode files to address add test + failures + +--- + stablehlo/dialect/Base.cpp | 18 ++++++++++-------- + stablehlo/dialect/StablehloOps.td | 6 +++--- + stablehlo/tests/ops_stablehlo.mlir | 12 ++++++------ + .../stablehlo_legalize_to_vhlo.0_10_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_10_0.mlir.bc | Bin 20200 -> 20200 bytes + .../stablehlo_legalize_to_vhlo.0_11_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_11_0.mlir.bc | Bin 20285 -> 20285 bytes + .../stablehlo_legalize_to_vhlo.0_12_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_12_0.mlir.bc | Bin 20929 -> 20929 bytes + .../stablehlo_legalize_to_vhlo.0_13_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_13_0.mlir.bc | Bin 20929 -> 20929 bytes + .../stablehlo_legalize_to_vhlo.0_14_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_14_0.mlir.bc | Bin 20579 -> 20579 bytes + .../stablehlo_legalize_to_vhlo.0_15_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_15_0.mlir.bc | Bin 16604 -> 16604 bytes + .../stablehlo_legalize_to_vhlo.0_16_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_16_0.mlir.bc | Bin 16732 -> 16732 bytes + .../stablehlo_legalize_to_vhlo.0_17_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_17_0.mlir.bc | Bin 17622 -> 17622 bytes + .../stablehlo_legalize_to_vhlo.0_18_0.mlir | 10 +++++----- + .../stablehlo_legalize_to_vhlo.0_18_0.mlir.bc | Bin 17750 -> 17739 bytes + .../stablehlo_legalize_to_vhlo.0_9_0.mlir | 8 ++++---- + .../stablehlo_legalize_to_vhlo.0_9_0.mlir.bc | Bin 20035 -> 20035 bytes + .../vhlo/stablehlo_legalize_to_vhlo.mlir | 8 ++++---- + 24 files changed, 64 insertions(+), 62 deletions(-) + +diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp +index 998dd183eb..299b1635f0 100644 +--- a/stablehlo/dialect/Base.cpp ++++ b/stablehlo/dialect/Base.cpp +@@ -70,11 +70,11 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { + // Q: Quantized (per-tensor or per-axis) + // NQ: Non Quantized + // tp1 tp2 Result +-// NQ NQ true/false ++// NQ NQ tp1 == tp2 + // NQ Q false +-// Q(per-tensor) Q(per-tensor) true/false ++// Q Q tp1.storage_type() == tp2.storage_type() || .. ++// Q(per-axis) Q(per-axis) tp1.quantized_dimension() == ... + // Q(per-tensor) Q(per-axis) false +-// Q(per-axis) Q(per-axis) true/false + + bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + // Get element type if shaped +@@ -83,13 +83,13 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + + auto qtp1 = tp1.dyn_cast(); + auto qtp2 = tp2.dyn_cast(); +- +- // If both are non quantized ++ + if (!qtp1 && !qtp2) { ++ // Both are non quantized + return tp1 == tp2; + } +- // If both are per-tensor Quantized + if (qtp1 && qtp2) { ++ // Both are quantized + if (qtp1.getStorageType() != qtp2.getStorageType() || + qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || + qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax() || +@@ -99,13 +99,15 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + + auto qpatp1 = qtp1.dyn_cast(); + auto qpatp2 = qtp2.dyn_cast(); +- // If both are also per-axis quantized + if(qpatp1 && qpatp2){ ++ // Both are also per-axis quantized + return qpatp1.getQuantizedDimension() == qpatp2.getQuantizedDimension(); + } ++ // return true if both are per-tensor quantized + return !(qpatp1 || qpatp2); ++ return true; + } +- ++ + return false; + } + +diff --git a/stablehlo/dialect/StablehloOps.td b/stablehlo/dialect/StablehloOps.td +index 29cdac7f95..295a6e6c3a 100644 +--- a/stablehlo/dialect/StablehloOps.td ++++ b/stablehlo/dialect/StablehloOps.td +@@ -1847,7 +1847,7 @@ def StableHLO_BroadcastOp : StableHLO_ShapedInterfaceOp<"broadcast", + } + + def StableHLO_BroadcastInDimOp : StableHLO_Op<"broadcast_in_dim", +- [Pure, /*can't use the Trait, need updates to verifierHLO_CompatibleOperandsAndResultElementType*/ /*broadcast_in_dim_c1*/]> { ++ [Pure, HLO_CompatibleOperandsAndResultElementType /*broadcast_in_dim_c1*/]> { + let summary = "BroadcastInDim operation"; + let description = [{ + Expands the dimensions and/or rank of an input tensor by duplicating the +@@ -2455,7 +2455,7 @@ def StableHLO_MapOp: StableHLO_ShapedInterfaceOp<"map", + } + + def StableHLO_ReshapeOp: StableHLO_Op<"reshape", +- [Pure, /* can't use the Trait, need verifier updatesHLO_CompatibleOperandsAndResultElementType*/]> { ++ [Pure, HLO_CompatibleOperandsAndResultElementType]> { + let summary = "Reshape operation"; + let description = [{ + Performs reshape of `operand` tensor to a `result` tensor. +@@ -2784,7 +2784,7 @@ def StableHLO_TraceOp: StableHLO_Op<"trace", []> { + } + + def StableHLO_TransposeOp: StableHLO_ShapedInterfaceOp<"transpose", +- [Pure, /* can't use the Trait, need verifier to handle per-axis HLO_CompatibleOperandsAndResultElementType,*/ ++ [Pure, HLO_CompatibleOperandsAndResultElementType, + DeclareOpInterfaceMethods]> { + let summary = "Transpose operation"; + let description = [{ +diff --git a/stablehlo/tests/ops_stablehlo.mlir b/stablehlo/tests/ops_stablehlo.mlir +index d5cabe763e..2d2ca1d014 100644 +--- a/stablehlo/tests/ops_stablehlo.mlir ++++ b/stablehlo/tests/ops_stablehlo.mlir +@@ -5396,11 +5396,11 @@ func.func @is_compatible_dynamism_dim_mismatch(%arg0: tensor<1x?xf32>) { + // TODO(b/230263270): For stablehlo.add, the plan is to only allow fp+fp=fp, q+q=q and q+q=fp. + func.func @is_compatible_quant_mix_non_quant(%arg0: tensor<1xf32>, %arg1: tensor<1x!quant.uniform>) { + %0 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32> +- %1 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1x!quant.uniform> +- %2 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> +- %3 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> +- %4 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> +- %5 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> ++ //%1 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1x!quant.uniform> ++ //%2 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> ++ //%3 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> ++ //%4 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> ++ //%5 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> + %6 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> + %7 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> + func.return +@@ -5474,7 +5474,7 @@ func.func @quantization_supported_ops(%arg0: tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>, %arg1: tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) { +- %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> ++ // %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> + %1 = "stablehlo.broadcast_in_dim"(%arg1) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) -> tensor<2x2x2x!quant.uniform:f32:0, {0.1:-30, 0.1:-30}>> + // %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> + // %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir +index d2eef1d28e..47c3010cf6 100644 +--- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir ++++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir +@@ -2233,10 +2233,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir.bc +index c040c899c3199b90c7a68116f9bceb38916aa835..94dd3fba5d17280e18d2e7d91184e5008b12c348 100644 +GIT binary patch +delta 945 +zcmY+CO-K}B7{__`@%43fT8FnWZFhE^#&JJh$Ng$q%$6mTAnM>pMM4J!Ln=B%hYS@7 +z3llmR^imNa_Ca;%5==o|I&|p}bm-v0qFW#Y1`%}anGtklU|wdP|MP!-|9NI(1vXaT +z^*c07Zkl7R*}+}cvaGEC2@^c(IHn^J7fi?h6&-5ToxEfN16(32s-yAlCJ8qp>Oy?9 +z)scL#xeF&XstK+>$Lp*pQ7fugVgwvDo=EH^3ESuJ;Lf;YQghVp#4r`t6ONh2v_Qlu +zfm4Wr#RcY7tJU3G#i~HM1S?nYj$l`>;f!F}3Cs#!uFT=Q;C=nqaY3-Yel>1cduoeI +zSQA_-K&ck3lO06^4Z%DwuqBUs0#Fx(ziGAWS-~)cb-}T&jHEo{0jV_8Fbp)Q1PFqi +z2Yf_5ZpgPn&a31sn5hMZnuc1yT^H=hGkh)(AiyPjz>9REC(u^Q$^Brfm&tOI>!Pc^ +zCr6uX8*RUwS_3x?F%A)qX=-Vr`Aon^&XMUDR=(1ydhB7yt64kQnP(&+86(*m$+*tq +zAi&^of@9u)V_ycX3P(6R)v>C2onu&rikf!Lo@(sOw1k~mn1DVG_jSxR*!ECS(>e9Q +z98@{Xa-8QhF&N8fY+A3{+rXAhndLS^LPz{Ec%?Fe}OP45b{xj}rO +z_%Fx_wYdoW5~?Ntb0_zYXqdaJ?F%XB14>xkioYI3>6Z#NK;0iHwu9;IthzHAuG&^J +gr4-oygVGRaF1Wy($=}~BUjoerJAAbL$;vJ3KRykWdH?_b + +delta 944 +zcmY*XO=uHQ5N0xcech~Scy4V>HtD+A=I1rdpCQB;g0u*thyJKo@X$h~h$x5$D{U!K +zr4~F?@S!47>JRF{i?9lM@!-XS;K4%=RlErWp@Ilro!vs4@?eLZ`R1GNd-FC{U}FVd +zzr!#$j4{WEaK|xCGh?`)Fvf$HZCDay-mrQDwx*Z#);(E+B9CQRCP6?bs +z6f7z*r#j8v{t8wE(h;m!#yf&ty@oS_r6w>Xc(FW(^MY4v*Kt9xYE6xs=AQcE64nLJ +zD8P(bG*3qIVGIi<$%M~wPXOwI)Nh*IdP*=%U_)?hsBm1Kkw7TT2#3RHPzewOI}iAX +zdhC>^Z-tmwiCHj1^9VH!H80-~?8!5HE)c+zOY#9PP(&%Psg~pWp{ZWR%Wbxew)!3) +zZEF&@TTZNj<3NZ*fTKe#Z8V<=_{cdj9sSBzimJ~VhMbzU;ypP=5P~s+tr3jtED8d` +z98Pe|xo@n?pjF@qho?GLRKI-;i%?e6_SsXdo#_t0GX)b+<8WWcR10k%6*ZkxAIw3O +z!z{;nP7}ScoW>Tm>Q$&ICs{x*o=*M_)nO5qz=^=Dj&(KS$WR6p0N4?j)@c`|!XfvA +zlMigk8}j7+O+AFXT29AMf>f{4$9=K?(i4;&%S=~VC2#JMd~mnqKiiV;W)r@m1KG)* +zR#WYo?-(Hg?>Y48wD8B-E@-H?*%V#vx9p&{y$F?-hdmPaoU{TU-J7{9&}Dn}f$YB^ +zCe-Ev^dwYA{^ySGA3@)Dvg->W=mImax@CVoh#9w3umBqFNU<9XxwGn?pug&-nkglq +e={_h8fo6jZyq);n&GJRin#AJ$f;(BcW&Q{0=ak|A + +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir +index 59657615db..d7d51de1b4 100644 +--- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir ++++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir +@@ -2240,10 +2240,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir.bc +index 2e3fc977ed958537bc528c8e4d7fd369d70685ad..314cc89d59f01107e5425936afc4fd319169ce1f 100644 +GIT binary patch +delta 933 +zcmY*XO-NKx6z0sC>+Q}fB5W#G@9UD>Lub6G5FvSsZhT{w( +zCsB@RnKnDR3s~T&7ALtpu5$9=AcUPK*5uxX4MdUsg=>Y-xD~OPq=| +zKq23lV|`g2bxu4lpj8}CIiSXgI$h_eqro);bUpYW7_bsdV3|`a%S>Dhc%agl@&5uK +zaB>&$g(7ikf;UyEEmevWg9Q~BLlzh*%hh;$*yCu)xAAtcj?tZ0Kjg8#~KQM$GH=#E13wC#5&i{XF-R920;BpvAwzy4?KC$ +zIb0K)a3>)vEAH%#laha&l6>*BKD0OiIVkd7m5XdHq$V#j=`Y6( +z`WC4*NHhxAs(0ovqTb$)(5MzHDC@nKvmNm +TzomBxm<3kt*YaQIpPByv<+7J3 + +delta 939 +zcmY*XO=uHA6z1h+I!@MfnQm>GY~s4vrs+tS9 +z>eMjC9V0>=$285f;q74;4OzBf3B(1%@_vSHwrZ9u$j}fKNL)D??UGB;Yi*XG!!1di +zW-868fW?43hz=1~AaiAJY!ofQM!}t`Rz~ynJp?6f#7iC?ukgvZp9gjig`-N)%pcm^xPMS4E~#zfU9+ScHx$%M}&*1dMt1s(QV0QFi^U&_UvX-LVf +zo_I&)7>zd8(2P)#pL$YiogjwUp#}1a81(gf%fbqP@U}AdpxyStL4oh6USx9t<-E+g +zUyc^^E!0?$2&ZAQxtW8I`uDb-0r@2R6)N6rt^*8lWN^3gTNzxH8+{3=diB1mz-+LA +Ws-@Xpd;cP^ssi#q!FyeJX8r@O@t0Nr + +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir +index 8d3f1a1405..46b5bf29b7 100644 +--- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir ++++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir +@@ -2240,10 +2240,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir.bc +index 684a242efbb87b704c0d5de38aa07f84fb2c57b0..a4e3a88fb6373af03bf33f77902390ba33a23e1e 100644 +GIT binary patch +delta 942 +zcmZ8fPe>F|9Nypl=J}kN+T}Tz?#^!G?Cdta9e3AtT~}MTQ0UUdl0b};B%HQP39j<{gh{!i^jM7MLC3=L6^q#$m_(h7&Auvm;;jo6$H +zH>Ge&l^CO9KQ=@P9Myv8ieCX+%*W#`@(5|QX;2B}gsZlA6(5n{uefbwFvAdWn&C7e +zBQb_Kku%d>Wh^t44VluiQ5$j;BZeAGP*qd|BQ>#_XjeRjmUxqB2TOcR6cxuoNBl{QE0%?p +zKbC9`yBP?brl3*GAfO_r5eB=IDjpIY55YaT5KGoLEJ8^bsayRyBAFx`k!<;9suYpx#Dz%+RVX~y(QUwcBwj&Pr$O`ER0mdIKzvH| +z25(<=G>BAT8CD?Vw84^&Rbi*|5e{nr5VG6gp^n_2NWTmpX}KO~iMX$!Ahz8kdfK$M5IgKiAaLJa~>9N7`zxK?4=70ks +W+9ap=Q@!s&v%ot1w*0ewo916)s+X?- + +delta 948 +zcmZ8fPe>F|7~gNdeLiOyrlthFXL_rqKNjmMNgU4=N>=1Yex<-Ev_TFl^#lSajzVCg%@AvzCQ`?5xHoSe0 +z5o($ft{J7SYgv|O`g@q7y6u=YM_e*(|A&4(Y}h$YMn)({k`p&#DX|*6(rR-$+LXi@ +zO{N*;uo4spu@REvsO3a=$7Qs+Ce~wSRQ0vU+oa8*DiX>G*KF~!V+?HZv%@i6bQvPf +zGMq(ZB*ri+a#pIlh((6-VWg*k>x|6b!v#jt)0k$or?80k867As;W8rwC9!B(r^>5q +zSZ37Q0QAIjYkVjd!3ZNzzU;S5&1Yj$3eDf +z=HCl|z{oAYClbU-E5c9Zs4Y2)5mQqWGz~SCuZZ<{dx$fvir4XWsEQBqLdbE@5x?V; +zp{fM(Clbv;&xJq>1)ZXcfFU`JFxaCsltp6TQE^`$#ELx$S@GGnWjs{Tw>uw!UWC8} +z3OfcCMa>z*9w-Pid2=XBB$7lU5-s0MLj|M;Ny9V*N)(9e5OAXKz(CKRO1%gkbGfFtguPdg6T4ml +zIdJjG8?Q6VciooR&93Et_{1vsu+kv1 +zRo072Btk`487^JCn7O2%xGy#{qjkmmnW-Rfe>nFTPCLp&kRI;Hb_-Ek2+F;SZ~4Op +zPrmAPPtwBf)ub<5$G#ls^8cGk4f}U{g8&tOu26zDezN4VIt_Vh$yHJ7bD-*f>6-)H +Z0S5xKNlx)+`rm) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir.bc +index 847748925eb47be3dcb888d980e0c770b7fb3c8f..2e15e9c5e661d3a271416ed0d43f260473cadd13 100644 +GIT binary patch +delta 942 +zcmZ8fPe>F|9Nypl=J}kN+T}Tz?#^!G?Cdta9e3AtT~}MTQ0UUdl0b};B%HQP39j<{gh{!i^jM7MLC3=L6^q#$m_(h7&Auvm;;jo6$H +zH>Ge&l^CO9KQ=@P9Myv8ieCX+%*W#`@(5|QX;2B}gsZlA6(5n{uefbwFvAdWn&C7e +zBQb_Kku%d>Wh^t44VluiQ5$j;BZeAGP*qd|BQ>#_XjeRjmUxqB2TOcR6cxuoNBl{QE0%?p +zKbC9`yBP?brl3*GAfO_r5eB=IDjpIY55YaT5KGoLEJ8^bsayRyBAFx`k!<;9suYpx#Dz%+RVX~y(QUwcBwj&Pr$O`ER0mdIKzvH| +z25(<=G>BAT8CD?Vw84^&Rbi*|5e{nr5VG6gp^n_2NWTmpX}KO~iMX$!Ahz8kdfK$M5IgKiAaLJa~>9N7`zxK?4=70ks +W+9ap=Q@!s&v%ot1w*0ewo916)s+X?- + +delta 948 +zcmZ8fPe>F|7~gNdeLiOyrlthFXL_rqKNjmMNgU4=N>=1Yex<-Ev_TFl^#lSajzVCg%@AvzCQ`?5xHoSe0 +z5o($ft{J7SYgv|O`g@q7y6u=YM_e*(|A&4(Y}h$YMn)({k`p&#DX|*6(rR-$+LXi@ +zO{N*;uo4spu@REvsO3a=$7Qs+Ce~wSRQ0vU+oa8*DiX>G*KF~!V+?HZv%@i6bQvPf +zGMq(ZB*ri+a#pIlh((6-VWg*k>x|6b!v#jt)0k$or?80k867As;W8rwC9!B(r^>5q +zSZ37Q0QAIjYkVjd!3ZNzzU;S5&1Yj$3eDf +z=HCl|z{oAYClbU-E5c9Zs4Y2)5mQqWGz~SCuZZ<{dx$fvir4XWsEQBqLdbE@5x?V; +zp{fM(Clbv;&xJq>1)ZXcfFU`JFxaCsltp6TQE^`$#ELx$S@GGnWjs{Tw>uw!UWC8} +z3OfcCMa>z*9w-Pid2=XBB$7lU5-s0MLj|M;Ny9V*N)(9e5OAXKz(CKRO1%gkbGfFtguPdg6T4ml +zIdJjG8?Q6VciooR&93Et_{1vsu+kv1 +zRo072Btk`487^JCn7O2%xGy#{qjkmmnW-Rfe>nFTPCLp&kRI;Hb_-Ek2+F;SZ~4Op +zPrmAPPtwBf)ub<5$G#ls^8cGk4f}U{g8&tOu26zDezN4VIt_Vh$yHJ7bD-*f>6-)H +Z0S5xKNlx)+`rm) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir.bc +index 4b48cf0681aa84aedcc60e16d351358b40e92e13..acb4b340f6eb1ed8bbbbace194704a143edf0bc3 100644 +GIT binary patch +delta 922 +zcmY+CPe>F|9LM)Jzj;1qN8Mh7Zo9MVFgvc}b6wX}7ZcgaLMSLq$&$z=4H1zJ9U@Fp +z)Ud=;JoOQ-N66m;ki5hNnLSs|5y-{bLpzwi6}zHc^`U}Fhh +zzQT|}4da4igsJ10rs;Ygv4#e1+psy}v|)RHgO@6FJIzUFC*?>o;&P-$WgjCwpuY_{KHCsH7c1v+HY8y#RGDI9< +zID*JXgkf4Vnl%kM%rTT7BiRQQ%PpK@w5@#_XBla07n7!0)^TSZ +zI~e8n0l8Q)d&8L!h8Rhud|T=9kOAsU;#b6~be9o~W1dmWi%g84Pk}t88KF=J4JiQv +zBcp(evd1a%J(kYP(wPxM^9?l(HJ>}r$kP>k#t^{AmH2>Ta>W!wORUBYfhFF@vZ~{t +zBer7ws%4?&^~Zlg@>HOlLXcuowy5%c!r%j?s*G62qPDJMQjA&WAT3^7aaq^>>TOxG +zpp6EOQdre7C+^y(mBC^0#V+L1MA}FzBCQ%rshUQrlQVAuZ4{p9=T(ZWoq|LYlBox +m`fr*Si&+aSZ!KGZ(yp!mmDxC+)A9&P^ZDY=^4_*Cng0M{uBQ|L + +delta 921 +zcmY+CUq}>D6vp?=+3R(7)a^FtwmZ8Hv*Wm4*Ku7}FA5ejjVtm|`Lp^RJOgi~5mJPPl4{=do^?+>BYejW$EXVTQwq +zj6@lxMWa#Eki#59`7x5o;yp%gjo~;W&Hy@$X0ms1lF`<-DV$-XwM|SIMp^sad2DAi +zbssP#R*c?=8^$mrTsEwfgM%o$LJ +z0>>$=YM2vq)>-AsfcRn+Thm0!NGT$v8cS)=MXHfAZvt%;o@nUo%Qnkw7FCVP%8_KX +zGICwaC!76cExE;R#FV5Gup=<5VLO$2;!sOOG`YbjmC8j~yWrP$^(qjsA~3CC%G
    +j>6lo|n$Y2`Ws6W+C) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir.bc +index 6f6da967ff4d7fc540769f5059f00a716e0969b2..795ccbbdefa222cb819917799d11c2a53477b976 100644 +GIT binary patch +delta 934 +zcmY*XPe>F|7@u#xc|ONk8J>fGw$-(9T$XiPw#AlJWZ4FBF{1Dfnir!DWf3VTuuw@; +zi^#gvL;6EzA`%fqFxM`fJanioo;(yCi}2DR579Sk);7cNX6Ajr@Av0>I}5P00I%oN +zfOW>P>bc{DLZPU)tnO4mQWo{7x~$5(`Z>9)=U6g^dQu*%+O;AW5r|k6SVR;oBrqW} +zM(bc2(*o%hEY*(p1iLzovw}qn7!^F#ehudZ&vY)}l3Eg1wao!;mw<<88G(gi6;fec4{Ib*k|F|7@u#xdA=QIWq1y{+UnZ4F3Y+t+hWTq%d!pPVuaxzGCe447>h_jfrUz% +zT10lrhxCWcL?j}JV6I&{dFW7GJb5TO7Sg3d9ing6tZjzj&Aj*h{(Qf0_7-7p5nj!! +z{&Qg~z{BBSFc`Jm6?Lx)60)R6<*GieuBdWfcUU5ZdP1J4{$_<~z|V96BT6YSay&IuMNVnpy{>N+k6p3W@dvS8^9-V$gqFa$klcyz=x2`B~c +zx_t+`1dpx(M&(OmurpzzDOfD#l@RlN0fb<*`%s{XmM|l1n|0|E&-=##hAbxLJFUA +z<=|Sn;rzm%R0Eh3ET6}GonVfGj$qyLqu|PkJycC(JX=aYlz!z^Gt|YjNZTOWTTBfU^|(J2MDz_ltQ6d>52_0q$al +zEQ(%rtSDcIk^V~X#vkyhH%g`en;5Ae-r`ykb%T!;e<^9d@-s$8; +zIaHfeX}MXu@4tLRZB>=) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir.bc +index b9b4c99c6a728ce95d8385545a135d4c6bcae444..d2ec74d21984b8af366430110bb9e1c494c727ad 100644 +GIT binary patch +delta 877 +zcmY*XO-vI(6y7&)hv~Adb-IyKmV(=*X=-hnLK+~XfofIO!lBhJ##Q#6|Z{IV2l?I#+d?AH<|40y!ZXg_jaueYh`#{CVuO< +zYlWHX+P0nW?vafe$cP0aM(z-?Vw@C#+HRUjqLC5(wd+ +zZy0hbju$2C3(3lX#c32A(p@4DOqXQ!CgM&;@AT_a4>BYFX6+9MAQ(2S^ +zR5lDUJ%!Aqxu_N_LP_TNGiXEF3+%oKz6q_n2zN0j=J_cyc2<1jBdyior7iH$_gZ@q +zmN8L9{EF!rG)yL8Q=vw3VM^Q&Wz_3k=!gGu1BnsM9fnox5^GJa_+t6Y{s97-GYlUv +z>8)Fn#Hh%Jh}1t__q=J;(LUCI%x=egD`9$Cxi6K36Yf(MFNMcpRD2Fk=c%!xx_BGukyvjd +zQY29F3XStn=c=4R<~)f^w-JpQ0&;AhqQ3+GGI)rU5cd1~%bNi_L;3ryTOHy^^DoGI +xJ$v`Tp484viO%%68frnzInCgBi_Qt~_4aB^#^ZX_eFDBhq4MQ;%Pnsj{{mzto8=qkA6;5X?)1()e=r~8XdzZXQ}jcoaHRjeiP?7&9vh!jyeUM%NKQ*PBKjZ3eGz3 +z%ws2KvB!Whv8MNRBsp?U<8ik|GC$!U5mfFzlLa)mZUE=*f}*QMu^F?RVO9)!;}93K +z-i84y>Pjp~)|Zl%Q$usfGz~R(Kg;Rj5-xKDa1UV01zzB2Vw__RA&o0c*?P)1@_%qs +zssiT3cW+v?EVM+#H>~CW!I|-Wg-EyJD?lD?*`op??S&&ujUm;sQt3up-1eUcC#WpS +z1}Yo+nVLYR(@a?L7hq84`O|Mg!tw3B1fFxsp#t2)HZjLflT%~j8y{}11}|=dkG)sg +z3$TRID&l3Prcl$FgiQq+$;3tRAdqrD?*c!(Qw=0S6nhBPuv4rz+2V`g(R=y`DAo{s +zz?idPOcJdkA0SfyOx=s75li`412Vf^^R0xbN!NX`D6C+wYw>b$97e?FU@22=^x972 +zmeU#XpVhLgrnkM$*EhK)SL)9~RO|_bgV`#*R}~t0D0)R44n-j$&WDbAc2XB_LtPT< +zY=#O1icY?94(cp7XOKBhB2!I7ZJK}_+o$kv-@gnVVI_ov-k$PS0MAigzYU{99B=*w +yS*L6N0oWJcoheb78g++S60=q_I7!el&_L+z*FDM95AoTZjGjeh|GhMOG# + +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir +index 5b06671a35..5f5a4c8a71 100644 +--- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir ++++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir +@@ -2383,10 +2383,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir.bc +index 8f459cf76b62d1f070d0f279ab6943555b8338d2..8066e685e9032c25113740917eab9b7259813b80 100644 +GIT binary patch +delta 935 +zcmY+BT}TvB6vyY>Gq=05>$u*gRdOBU6OgfT_=68Y4F1PK-tkn`@83S7=sUE +zFf~qeXjG>m?S56ms#N{JJ1ZT}HV|bmjt@)dHi_@)nalFE$$$>#kifa{zYO)3zoTgGpxeZ4s>8&2t<>7~(jBBb-_mTJB)wH)z6ffVR{t!6J0YW_+nKA>;n4ItjkBN}&kXu~nQ}-A7KG7W1nQ +zZYnoFD4Dz7D!C#(%)$4pJ57GTnL*ocyI7sPoZ<%KEZdu8|BdnLRdW--YmhL{PN +z?wd8wh_-ZLgvc8&*W6n>U@3E}0FAP$n+a2s9&b;N*k0S|L5^WR@7K--k*B0M5;fSr*LjEZ9s3o_zHWT$VX-{Ny* +zo2=okHH!p#+@sMnskXfigY57XWU7g%og*NpHyQgGEWQ2-{(HWdt<#~~{aE)Bpy1wa +wXo67c-xSxYOtri*wu`xV3@rCY{1Es$Ius_`W4i~eJK$??UpiUt;`)W?UvT8H_y7O^ + +delta 935 +zcmY+BUr1Cz6vpqFbH~}c>#E~*{jULeP+qP_r$R=j%=8{EWgo*`{l8H$fNP&n% +zmWEmcjSoRDSyV4l!Y#^|$fq7eNHC#@gdTbfdI%wcnyV{rcVG^4&-b1A%{l8+us#I~ +zv&7vQ()AJ+3^Je3r;9nV22pVpyX5`>4;Mu%iL(9e#6VQksEZJ_zNuxZRQ++(Nfphg +zAQPpkMO({Hm^7JWa;sz}(zvK`)`|}28S<8h7tVvuxJkK8qZSvN&UZCZ&3VGqG$ZMk +zCQXw+G>w_8Ij0JHN)7_4t6( +z8xweoQ?nmU&g$x?aE7zSraO3-)5a!oA=gL41A`bA)4AOxaUFF|BN4lvG|zEB7njt3 +zbBkRjrx?T}XP6WfZWAKnsC%E&Fl>fl&cb5cUFA=5`sy_j2PjV5=-t< +zNQuwxm}(eki2S@`Ji*a}9*(0p%4rG^tiinRP<32!ry+?UX;%S}wn2)iA*3$c%Cule +z3>Td6$EZw70hNMNObsJ*(L|+^pN4kXjIRYgh*`f2E`j5Ul1jr2Hi*lGy=35`SSvie +zKim8`V?Mc}G^F7PnpuUPF*Syoi%HvA-A_iw#5Fx?hXwsdUi<1trK|&%G0Y@}Sn~L+ +z_eC#>wsoPG$Qy5#JS^`sl+_$SJ*(;g!qkY(+tnqSOWSOy+0uUK6Msra>$1_tcZAz) +zljJ|RX_<{@=;q4;Hr7hnxj;7mqn!e@MOjS4Qz8f0=JmjJao%e{Ox*S!aqRS4eDNNV +zHLP-9nn0J;=Z}$M!|pK14o@OeeTdo^0Xe<-z_0ww>!0K9^Tl$78``Z;6>k8N*1f7) +u@I-c}SgkVEu*WzdR)YaZSwDl_;Am}Cm~4+>bsG1) -> tensor<*xf32> { + func.return %0 : tensor<*xf32> + } + +-// CHECK-LABEL: "type_per_tensor_quantization" +-func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++// CHECK-LABEL: "type_quantization" ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK-LABEL: "type_per_axis_quantization" +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir.bc +index 7ac3ab342767d1c4847bf8e4f17b6a584728e0e2..ed1b4ed419907bdc2e9db743333827bcece829cb 100644 +GIT binary patch +delta 991 +zcmY*XOK1~O6wQ5ip06|0*gmH*H6hfr9mi>mVT`ecSTsh1B9;1SwQ3bbP!TEOMu~zJ +zs#vKa4@FT>5XBFOI0$~I3l}1O7UFN$g$rE>f(th;#5*;`I0JK;ch0@%oVoLH1YV86 +zz4I(=ZE!3rD%#paJRVoq*eCE*1^emw0hcv-Nf!EhP1Y7qrL{-k`Cqk}8nNp19I<$n +z1=UKdHTwI8g3F6VL64+xk>}KWqk}zTYKt+T#*9Tq&gGI|)N=~W=d+0|Wx>^TyO^LR +zU8e(G7p_=3r4@J3o9=Gk&R9YHFdfT9S0dt|#6d*K6B4s3Ro_0R3%exJEqP}iuSh<8 +z0#8bw+Jz}8I`gM+REovLGk8|=#YI)AkMs1Ny_i;)>bFL77Fv>LGQm7*zAgbxNv<2i +z)6JCR*oGx3u%x`O3mLUJJlAmSz{3G4(n^I}7nCG_^aRHw0t6n&6W}4Ln2~5Bcz8}! +zzt7Oo;&1##96(!r3wxU5prd9^T_p<=BN&lb!HVRzjkbSys=->4uwi(IE7v6 +z#knXdddOIC4xq +zuriFw7wfyx|Mo*@MjyPwv_IbTj_DJUF-Dq?qgP>%+T0vx0|V+{b5CI4W%DjNn>VA< +zS8HR42`wE&`|IO4!&FV~5kj+p@T9M`?0;jYCkbQ;V>=D +z8?x$IY(?GV4%D>x5{l!mkCzz?_?HqMo9+aufl?nwp~VsPBMfMxpId&K)x>v^ezWnW +zHsd}qwgvwB{f6X2fRg{BwF4rVf2*CP3C#|6yHbhv7O;J%eJj*;cZb9j8tRYQ7eEc$ +dmRjiSHlJ?n<7OlrwL5M9j`KR=4|Tm4e*mxew6XvI + +delta 1006 +zcmY*XQD_ub6y5voJm1bt*7%&P8+TdaZZeLObzP=w8kS|7PPbhOQb}qPT?2~2V1w8# +zriIEXR3R24X$g;(77Hy%q*@ShC}^cW`6#8aAELH@e)yrFAoRl@Kg2t67uOkhmwD&h +zcg~snR_5W~dARZ&i|uwSDreYqIp|14Y)EV%N}iVJsa)5<)-skQ +z(kuDUF#aU@+*f!;a`zCrQVb2B#RVxUBj4bn`L+6%wf!_b6xuqMGGy-^ZD>R +zYF?6nrX<(vVjE0Xa_mQ63iOp1b0M!LW7`bd4n53Dky9#`ebblx`VIU;B0%VYJOLh~ +zig}4u1P{Lzt^afCOaD1OBMx9yJ&t*rZKJI=t=%C@65|+`*uz;{>c=F1Xioo9U~PNk*E#Cw_)$HQAeD^}70^cQiEc +zckdzEdt+mBqSMAHCbWDQ?WK>?4Bd|0^Mn>d;kK_2+3Uu>ZWG87CJ%zEK2D|+e#a4) +zRt4w?ep3CIw82x$$@in}8L0KCw<%8WQL4tE9-K>iY=a%921>$MRuSUCB6Pg{)c20?bKBxxHz&=p6V+G&a +z3lN=yh*(2vD(}pikM>M(GagIW!{y*t=Vsjabo%K2&!#__Z5&mL6)QMXSr#t) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_quantization" +-func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir.bc +index 805452951125ee099c35ea3774a62fecaca96113..b8ebd6b5f8835df3a28981ba73e5c9d2828ce4b1 100644 +GIT binary patch +delta 940 +zcmZ9KT}TvB6vsJx_U>(VTID*n+s>@s#&Mm!jO)5C>$WXRXkcWJeLMsShLR+`ghYu+ +zf%sOC#En2y!lxjJKIW-J9(?erJG+o>44lKAbN_#S^FOn`0_!XA +z`~@}})G*FFMuIwyWm&HO79%uj+lDO==MCH6i1szWGvYk7H~5p^T{(! +zwxENN4xG@`j8XX@StL1uT26jQo(5ZT(`->!$N`%Ml~6&1X3LZ2h#D@Nwvom(L&R~0 +zx{Ug=rWow%-}4erScr!Vx&}-+_Lsp<`=NS +zs8|FPa>zQH$Ted#Bc2!7Dv$dNP-hDNnHASvMz9U5jAB)8SOPq`W0@7FhO(qgmX!%3 +zh8Acw4YdHV%E;;?Tw@3jkSS8YN!8I~$mK}tAaFUA%7+~X9l4r17Umqe|2?${P8wnq +z{zNdXBw;`#3|>--M913Xwd(Rq`v_!ZmZwx54yu2dUj?lQM=9LVu_(6=9mhNrq_8iX +z%M#_HoQQIY@1sa{;=(ADDBRZ3-7_@RP}8Z>!3-2BT%kBil`ardC2p_JOHh`Zb`J)f +zcKnW2U>+90Nx+1T75OyV(K=MWC}0r)>;#PIDEzWB8LAWZ2ZWE)OOTTTT_$=e(Xp-* +zHQpy(7uDOSd!vE+YXh~}Nd4_U>g`A)_0L`EPwqrv?+0(!eG%H=6@?ZZt9#Pesp}A` +zaw?NnLEp;^YX6oaKWB~v3HWirL*V)S-ZZp2K~kWSdZ_+rHT8l`; +z1T5FHc^6f`VA(k=t)5>{^)VL*!`oebQ?#~vF1vf%fcwL}r$BSS0UA->{#w2W8V6oG +Ka(}DulJyt8S(h>Z + +delta 954 +zcmZ9KPe>F|9LM?X&$q9cX_e>LZacF&jpH~z)^T08b=#ID*1*Uh!wx}$p(I(CkSH-J +z5N`!Zya+@kJOx4YXHL~65;~R;2pvj@{v3l29z6AC7t)P^_jvF3zF$7S?{C&uV0{Ii +zzrZG{>*pLjK^@05O-~QrVw9?uty>&%Ublh`wXadLd`>DA%8~dYmh4bX4nq-+mCIb% +zhAu{uhnp=<2cqpbp~zQ^a<~~050ZJ}b5wlsA$bNYVHrmFj2y6OxDv_plVxebh^~ZMML{LJFEO(E2-g?_gk+Kwa8mBb8J5Lx>L8TGR4UtG+h~i`)bWNg +zAOzo2o8UMQqwpt+j+8V2B4O~7QY0EyE3f4)zO)TPTKHBq-Yv}GN8nUQpRO(;`@)WL6oTXA1im4R0 +z)#`aDh)uf(!=1MOjul}Z7QjisgoZ`&)NgMYtUdH#5diE2jA`fvMQ1WnBkT_eAG=H7 +zi%O?~VWMN5C#$?qIxomu)w@wg{k4wTsHgt+AN6*$p8Dr5^(Sv4x3}T#nlC~tyrR&o +zVQEhqJ2f3bNlbOQGU$6<1IoYQh|gU|!vumjw<5rU{`54o*kMv2lX|F7L}e}la=LH2 +zf9yu|UadyNV;(HmvU#V6ncUl!fHtxyhzdQXF5gAFt) +Uy@R!E9+a{)3Xc$M^<6ao0<-j)6#xJL + +diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir +index 4f0822483f..49e87338c0 100644 +--- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir ++++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir +@@ -2388,10 +2388,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { + } + + // CHECK-LABEL: "type_per_tensor_quantization" +-func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { +- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 +- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor +- func.return %0 : tensor ++func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { ++ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> ++ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> ++ func.return %0 : tensor> + } + + // CHECK-LABEL: "type_per_axis_quantization" + +From a09e8f63d1a120d01fab2fd0bf879e8721f1d00e Mon Sep 17 00:00:00 2001 +From: Abhinav +Date: Fri, 23 Feb 2024 18:12:25 +0000 +Subject: [PATCH 5/5] Relax per-axis dimension check + +--- + stablehlo/dialect/Base.cpp | 13 ++------- + stablehlo/tests/ops_stablehlo.mlir | 45 ++---------------------------- + 2 files changed, 4 insertions(+), 54 deletions(-) + +diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp +index 299b1635f0..781d7c0a39 100644 +--- a/stablehlo/dialect/Base.cpp ++++ b/stablehlo/dialect/Base.cpp +@@ -67,15 +67,6 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { + return success(); + } + +-// Q: Quantized (per-tensor or per-axis) +-// NQ: Non Quantized +-// tp1 tp2 Result +-// NQ NQ tp1 == tp2 +-// NQ Q false +-// Q Q tp1.storage_type() == tp2.storage_type() || .. +-// Q(per-axis) Q(per-axis) tp1.quantized_dimension() == ... +-// Q(per-tensor) Q(per-axis) false +- + bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + // Get element type if shaped + tp1 = getElementTypeOrSelf(tp1); +@@ -101,11 +92,11 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { + auto qpatp2 = qtp2.dyn_cast(); + if(qpatp1 && qpatp2){ + // Both are also per-axis quantized +- return qpatp1.getQuantizedDimension() == qpatp2.getQuantizedDimension(); ++ // For now, don't match dimentions, Per OP verifier will do it. ++ return true; + } + // return true if both are per-tensor quantized + return !(qpatp1 || qpatp2); +- return true; + } + + return false; +diff --git a/stablehlo/tests/ops_stablehlo.mlir b/stablehlo/tests/ops_stablehlo.mlir +index 2d2ca1d014..3a0063e63e 100644 +--- a/stablehlo/tests/ops_stablehlo.mlir ++++ b/stablehlo/tests/ops_stablehlo.mlir +@@ -5393,16 +5393,10 @@ func.func @is_compatible_dynamism_dim_mismatch(%arg0: tensor<1x?xf32>) { + + // ----- + +-// TODO(b/230263270): For stablehlo.add, the plan is to only allow fp+fp=fp, q+q=q and q+q=fp. + func.func @is_compatible_quant_mix_non_quant(%arg0: tensor<1xf32>, %arg1: tensor<1x!quant.uniform>) { + %0 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32> +- //%1 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1x!quant.uniform> +- //%2 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> +- //%3 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> +- //%4 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> +- //%5 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> +- %6 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> +- %7 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> ++ %1 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> ++ %2 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> + func.return + } + +@@ -5444,42 +5438,7 @@ func.func @is_compatible_quant_signedness_mismatch(%arg0: tensor<1x!quant.unifor + func.return + } + +-// ----- +- +-// The following is the not the exhaustive list of ops supporting quantized +-// types. The list will be updated as part of adding verification support for +-// quantized ops. +-func.func @quantization_supported_ops(%arg0: tensor<1x2x2x!quant.uniform>, %arg1: tensor<1x2x2x!quant.uniform>, %arg2: tensor>) { +- %0 = "stablehlo.atan2"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %1 = "stablehlo.divide"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %2 = "stablehlo.power"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %3 = "stablehlo.remainder"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %4 = "stablehlo.subtract"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- +- %5 = "stablehlo.abs"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %6 = "stablehlo.cbrt"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %7 = "stablehlo.cosine"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %8 = "stablehlo.exponential"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %9 = "stablehlo.exponential_minus_one"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %10 = "stablehlo.log"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %11 = "stablehlo.log_plus_one"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %12 = "stablehlo.logistic"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %13 = "stablehlo.negate"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %14 = "stablehlo.rsqrt"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %15 = "stablehlo.sign"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %16 = "stablehlo.sine"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %17 = "stablehlo.sqrt"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- %18 = "stablehlo.tanh"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> +- func.return +-} + +-func.func @per_axis_quantized_ops(%arg0: tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>, %arg1: tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) { +- // %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> +- %1 = "stablehlo.broadcast_in_dim"(%arg1) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) -> tensor<2x2x2x!quant.uniform:f32:0, {0.1:-30, 0.1:-30}>> +- // %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> +- // %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> +- func.return +-} + + // ----- + From 8f2d194ced9b53e45ed8c746ad7dcf4f71406d18 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Fri, 23 Feb 2024 19:09:34 +0000 Subject: [PATCH 02/14] remove old patch --- t2.patch | 1105 ------------------------------------------------------ 1 file changed, 1105 deletions(-) delete mode 100644 t2.patch diff --git a/t2.patch b/t2.patch deleted file mode 100644 index 5437a25b1ff..00000000000 --- a/t2.patch +++ /dev/null @@ -1,1105 +0,0 @@ -From cb8d9384082235518fb6c750b73959503400c740 Mon Sep 17 00:00:00 2001 -From: Abhinav -Date: Wed, 21 Feb 2024 19:46:10 +0000 -Subject: [PATCH 1/5] T2 Trait verifier - ---- - stablehlo/dialect/Base.cpp | 11 ++++++++++- - stablehlo/dialect/StablehloOps.td | 10 +++++----- - stablehlo/tests/ops_stablehlo.mlir | 4 ++-- - 3 files changed, 17 insertions(+), 8 deletions(-) - -diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp -index 9eb4f2340e..88b421b2d9 100644 ---- a/stablehlo/dialect/Base.cpp -+++ b/stablehlo/dialect/Base.cpp -@@ -83,7 +83,16 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || - qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) - return false; -- } -+ -+ auto qpatp1 = tp1.dyn_cast(); -+ auto qpatp2 = tp2.dyn_cast(); -+ if (qpatp1 && qpatp2) { -+ if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) -+ return false; -+ } -+ } -+ //if((qtp1 == NULL && qtp2 != NULL) || (qtp1 != NULL && qtp2 == NULL)) return false; -+ - auto etp1 = getExpressedTypeOrSelf(tp1); - auto etp2 = getExpressedTypeOrSelf(tp2); - -diff --git a/stablehlo/dialect/StablehloOps.td b/stablehlo/dialect/StablehloOps.td -index b2ab6d6389..29cdac7f95 100644 ---- a/stablehlo/dialect/StablehloOps.td -+++ b/stablehlo/dialect/StablehloOps.td -@@ -1847,7 +1847,7 @@ def StableHLO_BroadcastOp : StableHLO_ShapedInterfaceOp<"broadcast", - } - - def StableHLO_BroadcastInDimOp : StableHLO_Op<"broadcast_in_dim", -- [Pure, HLO_CompatibleOperandsAndResultElementType /*broadcast_in_dim_c1*/]> { -+ [Pure, /*can't use the Trait, need updates to verifierHLO_CompatibleOperandsAndResultElementType*/ /*broadcast_in_dim_c1*/]> { - let summary = "BroadcastInDim operation"; - let description = [{ - Expands the dimensions and/or rank of an input tensor by duplicating the -@@ -1926,7 +1926,7 @@ def StableHLO_DynamicBroadcastInDimOp : StableHLO_ShapedInterfaceOp< - // directly. - - def StableHLO_CholeskyOp : StableHLO_Op<"cholesky", -- [Pure, SameOperandsAndResultElementType /*cholesky_c1*/, -+ [Pure, HLO_CompatibleOperandsAndResultElementType /*cholesky_c1*/, - InferTensorType /*cholesky_c1*/]> { - let summary = "Cholesky operation"; - let description = [{ -@@ -1953,7 +1953,7 @@ def StableHLO_CholeskyOp : StableHLO_Op<"cholesky", - } - - def StableHLO_ClampOp : StableHLO_ShapedInterfaceOp<"clamp", [Pure, -- SameOperandsAndResultElementType /* clamp_c3 */, HLO_BroadcastingElementwise, -+ HLO_CompatibleOperandsAndResultElementType /* clamp_c3 */, HLO_BroadcastingElementwise, - InferTensorType]> { - let summary = "Clamp operation"; - let description = [{ -@@ -2455,7 +2455,7 @@ def StableHLO_MapOp: StableHLO_ShapedInterfaceOp<"map", - } - - def StableHLO_ReshapeOp: StableHLO_Op<"reshape", -- [Pure, HLO_CompatibleOperandsAndResultElementType]> { -+ [Pure, /* can't use the Trait, need verifier updatesHLO_CompatibleOperandsAndResultElementType*/]> { - let summary = "Reshape operation"; - let description = [{ - Performs reshape of `operand` tensor to a `result` tensor. -@@ -2784,7 +2784,7 @@ def StableHLO_TraceOp: StableHLO_Op<"trace", []> { - } - - def StableHLO_TransposeOp: StableHLO_ShapedInterfaceOp<"transpose", -- [Pure, HLO_CompatibleOperandsAndResultElementType, -+ [Pure, /* can't use the Trait, need verifier to handle per-axis HLO_CompatibleOperandsAndResultElementType,*/ - DeclareOpInterfaceMethods]> { - let summary = "Transpose operation"; - let description = [{ -diff --git a/stablehlo/tests/ops_stablehlo.mlir b/stablehlo/tests/ops_stablehlo.mlir -index 998054478f..d5cabe763e 100644 ---- a/stablehlo/tests/ops_stablehlo.mlir -+++ b/stablehlo/tests/ops_stablehlo.mlir -@@ -5476,8 +5476,8 @@ func.func @quantization_supported_ops(%arg0: tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>, %arg1: tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) { - %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> - %1 = "stablehlo.broadcast_in_dim"(%arg1) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) -> tensor<2x2x2x!quant.uniform:f32:0, {0.1:-30, 0.1:-30}>> -- %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> -- %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> -+ // %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> -+ // %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> - func.return - } - - -From cc743f716ff9d41db7a6dca6ede0ab20e483c155 Mon Sep 17 00:00:00 2001 -From: Abhinav -Date: Wed, 21 Feb 2024 20:34:04 +0000 -Subject: [PATCH 2/5] reorg isCompatibleElementTypeForHloTypeInference - ---- - stablehlo/dialect/Base.cpp | 62 ++++++++++++++++++++------------------ - 1 file changed, 33 insertions(+), 29 deletions(-) - -diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp -index 88b421b2d9..0ebcc01cba 100644 ---- a/stablehlo/dialect/Base.cpp -+++ b/stablehlo/dialect/Base.cpp -@@ -66,7 +66,9 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { - } - return success(); - } -- -+// tp1 tp2 Result -+// Quantized Non Quantized false -+// .... - bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - // Get element type if shaped - tp1 = getElementTypeOrSelf(tp1); -@@ -78,34 +80,36 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - // Individual ops may introduce additional constraints. - auto qtp1 = tp1.dyn_cast(); - auto qtp2 = tp2.dyn_cast(); -- if (qtp1 && qtp2) { -- if (qtp1.getStorageType() != qtp2.getStorageType() || -- qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || -- qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) -- return false; -- -- auto qpatp1 = tp1.dyn_cast(); -- auto qpatp2 = tp2.dyn_cast(); -- if (qpatp1 && qpatp2) { -- if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) -- return false; -- } -- } -- //if((qtp1 == NULL && qtp2 != NULL) || (qtp1 != NULL && qtp2 == NULL)) return false; -- -- auto etp1 = getExpressedTypeOrSelf(tp1); -- auto etp2 = getExpressedTypeOrSelf(tp2); -- -- // Sparsity: In the most general case, we allow any combination of -- // sparsity/denseness across any combination of operands/results, as well as -- // differences in sparsity encodings for operands and results. -- // Individual ops may introduce additional constraints. -- // No additional code is needed to check this because of how sparsity is -- // currently implemented. -- -- // Default case: Unless dynamism, quantization and/or sparsity are involved, -- // the types are required to be exactly equal. -- return etp1 == etp2; -+ if(!(qtp1 && qtp2)){ -+ // one Q and another non Q -+ if (qtp1 || qtp2) return false; -+ -+ // Sparsity: In the most general case, we allow any combination of -+ // sparsity/denseness across any combination of operands/results, as well as -+ // differences in sparsity encodings for operands and results. -+ // Individual ops may introduce additional constraints. -+ // No additional code is needed to check this because of how sparsity is -+ // currently implemented. -+ -+ // Default case: Unless dynamism, quantization and/or sparsity are involved, -+ // the types are required to be exactly equal. -+ -+ auto etp1 = getExpressedTypeOrSelf(tp1); -+ auto etp2 = getExpressedTypeOrSelf(tp2); -+ return etp1 == etp2; -+ } -+ if (qtp1.getStorageType() != qtp2.getStorageType() || -+ qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || -+ qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) -+ return false; -+ -+ auto qpatp1 = tp1.dyn_cast(); -+ auto qpatp2 = tp2.dyn_cast(); -+ if (qpatp1 && qpatp2) { -+ if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) -+ return false; -+ } -+ return true; - } - - bool isCompatibleForHloTypeInference(Type tp1, Type tp2) { - -From ac17026d5de29f30d94931645c3efe7d28cd506c Mon Sep 17 00:00:00 2001 -From: Abhinav -Date: Thu, 22 Feb 2024 18:58:13 +0000 -Subject: [PATCH 3/5] updated isCompatibleElementTypeForHloTypeInference logic - ---- - stablehlo/dialect/Base.cpp | 67 ++++++++++++++++++-------------------- - 1 file changed, 32 insertions(+), 35 deletions(-) - -diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp -index 0ebcc01cba..998dd183eb 100644 ---- a/stablehlo/dialect/Base.cpp -+++ b/stablehlo/dialect/Base.cpp -@@ -66,50 +66,47 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { - } - return success(); - } --// tp1 tp2 Result --// Quantized Non Quantized false --// .... -+ -+// Q: Quantized (per-tensor or per-axis) -+// NQ: Non Quantized -+// tp1 tp2 Result -+// NQ NQ true/false -+// NQ Q false -+// Q(per-tensor) Q(per-tensor) true/false -+// Q(per-tensor) Q(per-axis) false -+// Q(per-axis) Q(per-axis) true/false -+ - bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - // Get element type if shaped - tp1 = getElementTypeOrSelf(tp1); - tp2 = getElementTypeOrSelf(tp2); - -- // Quantization: In the most general case, we allow any combination of -- // quantized/non-quantized across any combination of operands/results, -- // and some differences in quantization parameters across operands/results. -- // Individual ops may introduce additional constraints. - auto qtp1 = tp1.dyn_cast(); - auto qtp2 = tp2.dyn_cast(); -- if(!(qtp1 && qtp2)){ -- // one Q and another non Q -- if (qtp1 || qtp2) return false; -- -- // Sparsity: In the most general case, we allow any combination of -- // sparsity/denseness across any combination of operands/results, as well as -- // differences in sparsity encodings for operands and results. -- // Individual ops may introduce additional constraints. -- // No additional code is needed to check this because of how sparsity is -- // currently implemented. -- -- // Default case: Unless dynamism, quantization and/or sparsity are involved, -- // the types are required to be exactly equal. -- -- auto etp1 = getExpressedTypeOrSelf(tp1); -- auto etp2 = getExpressedTypeOrSelf(tp2); -- return etp1 == etp2; -+ -+ // If both are non quantized -+ if (!qtp1 && !qtp2) { -+ return tp1 == tp2; - } -- if (qtp1.getStorageType() != qtp2.getStorageType() || -- qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || -- qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax()) -- return false; -- -- auto qpatp1 = tp1.dyn_cast(); -- auto qpatp2 = tp2.dyn_cast(); -- if (qpatp1 && qpatp2) { -- if(qpatp1.getQuantizedDimension() != qpatp2.getQuantizedDimension()) -- return false; -+ // If both are per-tensor Quantized -+ if (qtp1 && qtp2) { -+ if (qtp1.getStorageType() != qtp2.getStorageType() || -+ qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || -+ qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax() || -+ qtp1.getExpressedType() != qtp2.getExpressedType()){ -+ return false; -+ } -+ -+ auto qpatp1 = qtp1.dyn_cast(); -+ auto qpatp2 = qtp2.dyn_cast(); -+ // If both are also per-axis quantized -+ if(qpatp1 && qpatp2){ -+ return qpatp1.getQuantizedDimension() == qpatp2.getQuantizedDimension(); -+ } -+ return !(qpatp1 || qpatp2); - } -- return true; -+ -+ return false; - } - - bool isCompatibleForHloTypeInference(Type tp1, Type tp2) { - -From 7c342ce16fda44dc20313b05dd47120a0b860e39 Mon Sep 17 00:00:00 2001 -From: Abhinav -Date: Fri, 23 Feb 2024 00:38:54 +0000 -Subject: [PATCH 4/5] generated new vhlo bytecode files to address add test - failures - ---- - stablehlo/dialect/Base.cpp | 18 ++++++++++-------- - stablehlo/dialect/StablehloOps.td | 6 +++--- - stablehlo/tests/ops_stablehlo.mlir | 12 ++++++------ - .../stablehlo_legalize_to_vhlo.0_10_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_10_0.mlir.bc | Bin 20200 -> 20200 bytes - .../stablehlo_legalize_to_vhlo.0_11_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_11_0.mlir.bc | Bin 20285 -> 20285 bytes - .../stablehlo_legalize_to_vhlo.0_12_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_12_0.mlir.bc | Bin 20929 -> 20929 bytes - .../stablehlo_legalize_to_vhlo.0_13_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_13_0.mlir.bc | Bin 20929 -> 20929 bytes - .../stablehlo_legalize_to_vhlo.0_14_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_14_0.mlir.bc | Bin 20579 -> 20579 bytes - .../stablehlo_legalize_to_vhlo.0_15_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_15_0.mlir.bc | Bin 16604 -> 16604 bytes - .../stablehlo_legalize_to_vhlo.0_16_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_16_0.mlir.bc | Bin 16732 -> 16732 bytes - .../stablehlo_legalize_to_vhlo.0_17_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_17_0.mlir.bc | Bin 17622 -> 17622 bytes - .../stablehlo_legalize_to_vhlo.0_18_0.mlir | 10 +++++----- - .../stablehlo_legalize_to_vhlo.0_18_0.mlir.bc | Bin 17750 -> 17739 bytes - .../stablehlo_legalize_to_vhlo.0_9_0.mlir | 8 ++++---- - .../stablehlo_legalize_to_vhlo.0_9_0.mlir.bc | Bin 20035 -> 20035 bytes - .../vhlo/stablehlo_legalize_to_vhlo.mlir | 8 ++++---- - 24 files changed, 64 insertions(+), 62 deletions(-) - -diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp -index 998dd183eb..299b1635f0 100644 ---- a/stablehlo/dialect/Base.cpp -+++ b/stablehlo/dialect/Base.cpp -@@ -70,11 +70,11 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { - // Q: Quantized (per-tensor or per-axis) - // NQ: Non Quantized - // tp1 tp2 Result --// NQ NQ true/false -+// NQ NQ tp1 == tp2 - // NQ Q false --// Q(per-tensor) Q(per-tensor) true/false -+// Q Q tp1.storage_type() == tp2.storage_type() || .. -+// Q(per-axis) Q(per-axis) tp1.quantized_dimension() == ... - // Q(per-tensor) Q(per-axis) false --// Q(per-axis) Q(per-axis) true/false - - bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - // Get element type if shaped -@@ -83,13 +83,13 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - - auto qtp1 = tp1.dyn_cast(); - auto qtp2 = tp2.dyn_cast(); -- -- // If both are non quantized -+ - if (!qtp1 && !qtp2) { -+ // Both are non quantized - return tp1 == tp2; - } -- // If both are per-tensor Quantized - if (qtp1 && qtp2) { -+ // Both are quantized - if (qtp1.getStorageType() != qtp2.getStorageType() || - qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || - qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax() || -@@ -99,13 +99,15 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - - auto qpatp1 = qtp1.dyn_cast(); - auto qpatp2 = qtp2.dyn_cast(); -- // If both are also per-axis quantized - if(qpatp1 && qpatp2){ -+ // Both are also per-axis quantized - return qpatp1.getQuantizedDimension() == qpatp2.getQuantizedDimension(); - } -+ // return true if both are per-tensor quantized - return !(qpatp1 || qpatp2); -+ return true; - } -- -+ - return false; - } - -diff --git a/stablehlo/dialect/StablehloOps.td b/stablehlo/dialect/StablehloOps.td -index 29cdac7f95..295a6e6c3a 100644 ---- a/stablehlo/dialect/StablehloOps.td -+++ b/stablehlo/dialect/StablehloOps.td -@@ -1847,7 +1847,7 @@ def StableHLO_BroadcastOp : StableHLO_ShapedInterfaceOp<"broadcast", - } - - def StableHLO_BroadcastInDimOp : StableHLO_Op<"broadcast_in_dim", -- [Pure, /*can't use the Trait, need updates to verifierHLO_CompatibleOperandsAndResultElementType*/ /*broadcast_in_dim_c1*/]> { -+ [Pure, HLO_CompatibleOperandsAndResultElementType /*broadcast_in_dim_c1*/]> { - let summary = "BroadcastInDim operation"; - let description = [{ - Expands the dimensions and/or rank of an input tensor by duplicating the -@@ -2455,7 +2455,7 @@ def StableHLO_MapOp: StableHLO_ShapedInterfaceOp<"map", - } - - def StableHLO_ReshapeOp: StableHLO_Op<"reshape", -- [Pure, /* can't use the Trait, need verifier updatesHLO_CompatibleOperandsAndResultElementType*/]> { -+ [Pure, HLO_CompatibleOperandsAndResultElementType]> { - let summary = "Reshape operation"; - let description = [{ - Performs reshape of `operand` tensor to a `result` tensor. -@@ -2784,7 +2784,7 @@ def StableHLO_TraceOp: StableHLO_Op<"trace", []> { - } - - def StableHLO_TransposeOp: StableHLO_ShapedInterfaceOp<"transpose", -- [Pure, /* can't use the Trait, need verifier to handle per-axis HLO_CompatibleOperandsAndResultElementType,*/ -+ [Pure, HLO_CompatibleOperandsAndResultElementType, - DeclareOpInterfaceMethods]> { - let summary = "Transpose operation"; - let description = [{ -diff --git a/stablehlo/tests/ops_stablehlo.mlir b/stablehlo/tests/ops_stablehlo.mlir -index d5cabe763e..2d2ca1d014 100644 ---- a/stablehlo/tests/ops_stablehlo.mlir -+++ b/stablehlo/tests/ops_stablehlo.mlir -@@ -5396,11 +5396,11 @@ func.func @is_compatible_dynamism_dim_mismatch(%arg0: tensor<1x?xf32>) { - // TODO(b/230263270): For stablehlo.add, the plan is to only allow fp+fp=fp, q+q=q and q+q=fp. - func.func @is_compatible_quant_mix_non_quant(%arg0: tensor<1xf32>, %arg1: tensor<1x!quant.uniform>) { - %0 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32> -- %1 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1x!quant.uniform> -- %2 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -- %3 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -- %4 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> -- %5 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> -+ //%1 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1x!quant.uniform> -+ //%2 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -+ //%3 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -+ //%4 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> -+ //%5 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> - %6 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> - %7 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> - func.return -@@ -5474,7 +5474,7 @@ func.func @quantization_supported_ops(%arg0: tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>, %arg1: tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) { -- %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> -+ // %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> - %1 = "stablehlo.broadcast_in_dim"(%arg1) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) -> tensor<2x2x2x!quant.uniform:f32:0, {0.1:-30, 0.1:-30}>> - // %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> - // %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir -index d2eef1d28e..47c3010cf6 100644 ---- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir -+++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir -@@ -2233,10 +2233,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir.bc -index c040c899c3199b90c7a68116f9bceb38916aa835..94dd3fba5d17280e18d2e7d91184e5008b12c348 100644 -GIT binary patch -delta 945 -zcmY+CO-K}B7{__`@%43fT8FnWZFhE^#&JJh$Ng$q%$6mTAnM>pMM4J!Ln=B%hYS@7 -z3llmR^imNa_Ca;%5==o|I&|p}bm-v0qFW#Y1`%}anGtklU|wdP|MP!-|9NI(1vXaT -z^*c07Zkl7R*}+}cvaGEC2@^c(IHn^J7fi?h6&-5ToxEfN16(32s-yAlCJ8qp>Oy?9 -z)scL#xeF&XstK+>$Lp*pQ7fugVgwvDo=EH^3ESuJ;Lf;YQghVp#4r`t6ONh2v_Qlu -zfm4Wr#RcY7tJU3G#i~HM1S?nYj$l`>;f!F}3Cs#!uFT=Q;C=nqaY3-Yel>1cduoeI -zSQA_-K&ck3lO06^4Z%DwuqBUs0#Fx(ziGAWS-~)cb-}T&jHEo{0jV_8Fbp)Q1PFqi -z2Yf_5ZpgPn&a31sn5hMZnuc1yT^H=hGkh)(AiyPjz>9REC(u^Q$^Brfm&tOI>!Pc^ -zCr6uX8*RUwS_3x?F%A)qX=-Vr`Aon^&XMUDR=(1ydhB7yt64kQnP(&+86(*m$+*tq -zAi&^of@9u)V_ycX3P(6R)v>C2onu&rikf!Lo@(sOw1k~mn1DVG_jSxR*!ECS(>e9Q -z98@{Xa-8QhF&N8fY+A3{+rXAhndLS^LPz{Ec%?Fe}OP45b{xj}rO -z_%Fx_wYdoW5~?Ntb0_zYXqdaJ?F%XB14>xkioYI3>6Z#NK;0iHwu9;IthzHAuG&^J -gr4-oygVGRaF1Wy($=}~BUjoerJAAbL$;vJ3KRykWdH?_b - -delta 944 -zcmY*XO=uHQ5N0xcech~Scy4V>HtD+A=I1rdpCQB;g0u*thyJKo@X$h~h$x5$D{U!K -zr4~F?@S!47>JRF{i?9lM@!-XS;K4%=RlErWp@Ilro!vs4@?eLZ`R1GNd-FC{U}FVd -zzr!#$j4{WEaK|xCGh?`)Fvf$HZCDay-mrQDwx*Z#);(E+B9CQRCP6?bs -z6f7z*r#j8v{t8wE(h;m!#yf&ty@oS_r6w>Xc(FW(^MY4v*Kt9xYE6xs=AQcE64nLJ -zD8P(bG*3qIVGIi<$%M~wPXOwI)Nh*IdP*=%U_)?hsBm1Kkw7TT2#3RHPzewOI}iAX -zdhC>^Z-tmwiCHj1^9VH!H80-~?8!5HE)c+zOY#9PP(&%Psg~pWp{ZWR%Wbxew)!3) -zZEF&@TTZNj<3NZ*fTKe#Z8V<=_{cdj9sSBzimJ~VhMbzU;ypP=5P~s+tr3jtED8d` -z98Pe|xo@n?pjF@qho?GLRKI-;i%?e6_SsXdo#_t0GX)b+<8WWcR10k%6*ZkxAIw3O -z!z{;nP7}ScoW>Tm>Q$&ICs{x*o=*M_)nO5qz=^=Dj&(KS$WR6p0N4?j)@c`|!XfvA -zlMigk8}j7+O+AFXT29AMf>f{4$9=K?(i4;&%S=~VC2#JMd~mnqKiiV;W)r@m1KG)* -zR#WYo?-(Hg?>Y48wD8B-E@-H?*%V#vx9p&{y$F?-hdmPaoU{TU-J7{9&}Dn}f$YB^ -zCe-Ev^dwYA{^ySGA3@)Dvg->W=mImax@CVoh#9w3umBqFNU<9XxwGn?pug&-nkglq -e={_h8fo6jZyq);n&GJRin#AJ$f;(BcW&Q{0=ak|A - -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir -index 59657615db..d7d51de1b4 100644 ---- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir -+++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir -@@ -2240,10 +2240,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_11_0.mlir.bc -index 2e3fc977ed958537bc528c8e4d7fd369d70685ad..314cc89d59f01107e5425936afc4fd319169ce1f 100644 -GIT binary patch -delta 933 -zcmY*XO-NKx6z0sC>+Q}fB5W#G@9UD>Lub6G5FvSsZhT{w( -zCsB@RnKnDR3s~T&7ALtpu5$9=AcUPK*5uxX4MdUsg=>Y-xD~OPq=| -zKq23lV|`g2bxu4lpj8}CIiSXgI$h_eqro);bUpYW7_bsdV3|`a%S>Dhc%agl@&5uK -zaB>&$g(7ikf;UyEEmevWg9Q~BLlzh*%hh;$*yCu)xAAtcj?tZ0Kjg8#~KQM$GH=#E13wC#5&i{XF-R920;BpvAwzy4?KC$ -zIb0K)a3>)vEAH%#laha&l6>*BKD0OiIVkd7m5XdHq$V#j=`Y6( -z`WC4*NHhxAs(0ovqTb$)(5MzHDC@nKvmNm -TzomBxm<3kt*YaQIpPByv<+7J3 - -delta 939 -zcmY*XO=uHA6z1h+I!@MfnQm>GY~s4vrs+tS9 -z>eMjC9V0>=$285f;q74;4OzBf3B(1%@_vSHwrZ9u$j}fKNL)D??UGB;Yi*XG!!1di -zW-868fW?43hz=1~AaiAJY!ofQM!}t`Rz~ynJp?6f#7iC?ukgvZp9gjig`-N)%pcm^xPMS4E~#zfU9+ScHx$%M}&*1dMt1s(QV0QFi^U&_UvX-LVf -zo_I&)7>zd8(2P)#pL$YiogjwUp#}1a81(gf%fbqP@U}AdpxyStL4oh6USx9t<-E+g -zUyc^^E!0?$2&ZAQxtW8I`uDb-0r@2R6)N6rt^*8lWN^3gTNzxH8+{3=diB1mz-+LA -Ws-@Xpd;cP^ssi#q!FyeJX8r@O@t0Nr - -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir -index 8d3f1a1405..46b5bf29b7 100644 ---- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir -+++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir -@@ -2240,10 +2240,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir.bc -index 684a242efbb87b704c0d5de38aa07f84fb2c57b0..a4e3a88fb6373af03bf33f77902390ba33a23e1e 100644 -GIT binary patch -delta 942 -zcmZ8fPe>F|9Nypl=J}kN+T}Tz?#^!G?Cdta9e3AtT~}MTQ0UUdl0b};B%HQP39j<{gh{!i^jM7MLC3=L6^q#$m_(h7&Auvm;;jo6$H -zH>Ge&l^CO9KQ=@P9Myv8ieCX+%*W#`@(5|QX;2B}gsZlA6(5n{uefbwFvAdWn&C7e -zBQb_Kku%d>Wh^t44VluiQ5$j;BZeAGP*qd|BQ>#_XjeRjmUxqB2TOcR6cxuoNBl{QE0%?p -zKbC9`yBP?brl3*GAfO_r5eB=IDjpIY55YaT5KGoLEJ8^bsayRyBAFx`k!<;9suYpx#Dz%+RVX~y(QUwcBwj&Pr$O`ER0mdIKzvH| -z25(<=G>BAT8CD?Vw84^&Rbi*|5e{nr5VG6gp^n_2NWTmpX}KO~iMX$!Ahz8kdfK$M5IgKiAaLJa~>9N7`zxK?4=70ks -W+9ap=Q@!s&v%ot1w*0ewo916)s+X?- - -delta 948 -zcmZ8fPe>F|7~gNdeLiOyrlthFXL_rqKNjmMNgU4=N>=1Yex<-Ev_TFl^#lSajzVCg%@AvzCQ`?5xHoSe0 -z5o($ft{J7SYgv|O`g@q7y6u=YM_e*(|A&4(Y}h$YMn)({k`p&#DX|*6(rR-$+LXi@ -zO{N*;uo4spu@REvsO3a=$7Qs+Ce~wSRQ0vU+oa8*DiX>G*KF~!V+?HZv%@i6bQvPf -zGMq(ZB*ri+a#pIlh((6-VWg*k>x|6b!v#jt)0k$or?80k867As;W8rwC9!B(r^>5q -zSZ37Q0QAIjYkVjd!3ZNzzU;S5&1Yj$3eDf -z=HCl|z{oAYClbU-E5c9Zs4Y2)5mQqWGz~SCuZZ<{dx$fvir4XWsEQBqLdbE@5x?V; -zp{fM(Clbv;&xJq>1)ZXcfFU`JFxaCsltp6TQE^`$#ELx$S@GGnWjs{Tw>uw!UWC8} -z3OfcCMa>z*9w-Pid2=XBB$7lU5-s0MLj|M;Ny9V*N)(9e5OAXKz(CKRO1%gkbGfFtguPdg6T4ml -zIdJjG8?Q6VciooR&93Et_{1vsu+kv1 -zRo072Btk`487^JCn7O2%xGy#{qjkmmnW-Rfe>nFTPCLp&kRI;Hb_-Ek2+F;SZ~4Op -zPrmAPPtwBf)ub<5$G#ls^8cGk4f}U{g8&tOu26zDezN4VIt_Vh$yHJ7bD-*f>6-)H -Z0S5xKNlx)+`rm) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir.bc -index 847748925eb47be3dcb888d980e0c770b7fb3c8f..2e15e9c5e661d3a271416ed0d43f260473cadd13 100644 -GIT binary patch -delta 942 -zcmZ8fPe>F|9Nypl=J}kN+T}Tz?#^!G?Cdta9e3AtT~}MTQ0UUdl0b};B%HQP39j<{gh{!i^jM7MLC3=L6^q#$m_(h7&Auvm;;jo6$H -zH>Ge&l^CO9KQ=@P9Myv8ieCX+%*W#`@(5|QX;2B}gsZlA6(5n{uefbwFvAdWn&C7e -zBQb_Kku%d>Wh^t44VluiQ5$j;BZeAGP*qd|BQ>#_XjeRjmUxqB2TOcR6cxuoNBl{QE0%?p -zKbC9`yBP?brl3*GAfO_r5eB=IDjpIY55YaT5KGoLEJ8^bsayRyBAFx`k!<;9suYpx#Dz%+RVX~y(QUwcBwj&Pr$O`ER0mdIKzvH| -z25(<=G>BAT8CD?Vw84^&Rbi*|5e{nr5VG6gp^n_2NWTmpX}KO~iMX$!Ahz8kdfK$M5IgKiAaLJa~>9N7`zxK?4=70ks -W+9ap=Q@!s&v%ot1w*0ewo916)s+X?- - -delta 948 -zcmZ8fPe>F|7~gNdeLiOyrlthFXL_rqKNjmMNgU4=N>=1Yex<-Ev_TFl^#lSajzVCg%@AvzCQ`?5xHoSe0 -z5o($ft{J7SYgv|O`g@q7y6u=YM_e*(|A&4(Y}h$YMn)({k`p&#DX|*6(rR-$+LXi@ -zO{N*;uo4spu@REvsO3a=$7Qs+Ce~wSRQ0vU+oa8*DiX>G*KF~!V+?HZv%@i6bQvPf -zGMq(ZB*ri+a#pIlh((6-VWg*k>x|6b!v#jt)0k$or?80k867As;W8rwC9!B(r^>5q -zSZ37Q0QAIjYkVjd!3ZNzzU;S5&1Yj$3eDf -z=HCl|z{oAYClbU-E5c9Zs4Y2)5mQqWGz~SCuZZ<{dx$fvir4XWsEQBqLdbE@5x?V; -zp{fM(Clbv;&xJq>1)ZXcfFU`JFxaCsltp6TQE^`$#ELx$S@GGnWjs{Tw>uw!UWC8} -z3OfcCMa>z*9w-Pid2=XBB$7lU5-s0MLj|M;Ny9V*N)(9e5OAXKz(CKRO1%gkbGfFtguPdg6T4ml -zIdJjG8?Q6VciooR&93Et_{1vsu+kv1 -zRo072Btk`487^JCn7O2%xGy#{qjkmmnW-Rfe>nFTPCLp&kRI;Hb_-Ek2+F;SZ~4Op -zPrmAPPtwBf)ub<5$G#ls^8cGk4f}U{g8&tOu26zDezN4VIt_Vh$yHJ7bD-*f>6-)H -Z0S5xKNlx)+`rm) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir.bc -index 4b48cf0681aa84aedcc60e16d351358b40e92e13..acb4b340f6eb1ed8bbbbace194704a143edf0bc3 100644 -GIT binary patch -delta 922 -zcmY+CPe>F|9LM)Jzj;1qN8Mh7Zo9MVFgvc}b6wX}7ZcgaLMSLq$&$z=4H1zJ9U@Fp -z)Ud=;JoOQ-N66m;ki5hNnLSs|5y-{bLpzwi6}zHc^`U}Fhh -zzQT|}4da4igsJ10rs;Ygv4#e1+psy}v|)RHgO@6FJIzUFC*?>o;&P-$WgjCwpuY_{KHCsH7c1v+HY8y#RGDI9< -zID*JXgkf4Vnl%kM%rTT7BiRQQ%PpK@w5@#_XBla07n7!0)^TSZ -zI~e8n0l8Q)d&8L!h8Rhud|T=9kOAsU;#b6~be9o~W1dmWi%g84Pk}t88KF=J4JiQv -zBcp(evd1a%J(kYP(wPxM^9?l(HJ>}r$kP>k#t^{AmH2>Ta>W!wORUBYfhFF@vZ~{t -zBer7ws%4?&^~Zlg@>HOlLXcuowy5%c!r%j?s*G62qPDJMQjA&WAT3^7aaq^>>TOxG -zpp6EOQdre7C+^y(mBC^0#V+L1MA}FzBCQ%rshUQrlQVAuZ4{p9=T(ZWoq|LYlBox -m`fr*Si&+aSZ!KGZ(yp!mmDxC+)A9&P^ZDY=^4_*Cng0M{uBQ|L - -delta 921 -zcmY+CUq}>D6vp?=+3R(7)a^FtwmZ8Hv*Wm4*Ku7}FA5ejjVtm|`Lp^RJOgi~5mJPPl4{=do^?+>BYejW$EXVTQwq -zj6@lxMWa#Eki#59`7x5o;yp%gjo~;W&Hy@$X0ms1lF`<-DV$-XwM|SIMp^sad2DAi -zbssP#R*c?=8^$mrTsEwfgM%o$LJ -z0>>$=YM2vq)>-AsfcRn+Thm0!NGT$v8cS)=MXHfAZvt%;o@nUo%Qnkw7FCVP%8_KX -zGICwaC!76cExE;R#FV5Gup=<5VLO$2;!sOOG`YbjmC8j~yWrP$^(qjsA~3CC%G
      -j>6lo|n$Y2`Ws6W+C) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_15_0.mlir.bc -index 6f6da967ff4d7fc540769f5059f00a716e0969b2..795ccbbdefa222cb819917799d11c2a53477b976 100644 -GIT binary patch -delta 934 -zcmY*XPe>F|7@u#xc|ONk8J>fGw$-(9T$XiPw#AlJWZ4FBF{1Dfnir!DWf3VTuuw@; -zi^#gvL;6EzA`%fqFxM`fJanioo;(yCi}2DR579Sk);7cNX6Ajr@Av0>I}5P00I%oN -zfOW>P>bc{DLZPU)tnO4mQWo{7x~$5(`Z>9)=U6g^dQu*%+O;AW5r|k6SVR;oBrqW} -zM(bc2(*o%hEY*(p1iLzovw}qn7!^F#ehudZ&vY)}l3Eg1wao!;mw<<88G(gi6;fec4{Ib*k|F|7@u#xdA=QIWq1y{+UnZ4F3Y+t+hWTq%d!pPVuaxzGCe447>h_jfrUz% -zT10lrhxCWcL?j}JV6I&{dFW7GJb5TO7Sg3d9ing6tZjzj&Aj*h{(Qf0_7-7p5nj!! -z{&Qg~z{BBSFc`Jm6?Lx)60)R6<*GieuBdWfcUU5ZdP1J4{$_<~z|V96BT6YSay&IuMNVnpy{>N+k6p3W@dvS8^9-V$gqFa$klcyz=x2`B~c -zx_t+`1dpx(M&(OmurpzzDOfD#l@RlN0fb<*`%s{XmM|l1n|0|E&-=##hAbxLJFUA -z<=|Sn;rzm%R0Eh3ET6}GonVfGj$qyLqu|PkJycC(JX=aYlz!z^Gt|YjNZTOWTTBfU^|(J2MDz_ltQ6d>52_0q$al -zEQ(%rtSDcIk^V~X#vkyhH%g`en;5Ae-r`ykb%T!;e<^9d@-s$8; -zIaHfeX}MXu@4tLRZB>=) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir.bc -index b9b4c99c6a728ce95d8385545a135d4c6bcae444..d2ec74d21984b8af366430110bb9e1c494c727ad 100644 -GIT binary patch -delta 877 -zcmY*XO-vI(6y7&)hv~Adb-IyKmV(=*X=-hnLK+~XfofIO!lBhJ##Q#6|Z{IV2l?I#+d?AH<|40y!ZXg_jaueYh`#{CVuO< -zYlWHX+P0nW?vafe$cP0aM(z-?Vw@C#+HRUjqLC5(wd+ -zZy0hbju$2C3(3lX#c32A(p@4DOqXQ!CgM&;@AT_a4>BYFX6+9MAQ(2S^ -zR5lDUJ%!Aqxu_N_LP_TNGiXEF3+%oKz6q_n2zN0j=J_cyc2<1jBdyior7iH$_gZ@q -zmN8L9{EF!rG)yL8Q=vw3VM^Q&Wz_3k=!gGu1BnsM9fnox5^GJa_+t6Y{s97-GYlUv -z>8)Fn#Hh%Jh}1t__q=J;(LUCI%x=egD`9$Cxi6K36Yf(MFNMcpRD2Fk=c%!xx_BGukyvjd -zQY29F3XStn=c=4R<~)f^w-JpQ0&;AhqQ3+GGI)rU5cd1~%bNi_L;3ryTOHy^^DoGI -xJ$v`Tp484viO%%68frnzInCgBi_Qt~_4aB^#^ZX_eFDBhq4MQ;%Pnsj{{mzto8=qkA6;5X?)1()e=r~8XdzZXQ}jcoaHRjeiP?7&9vh!jyeUM%NKQ*PBKjZ3eGz3 -z%ws2KvB!Whv8MNRBsp?U<8ik|GC$!U5mfFzlLa)mZUE=*f}*QMu^F?RVO9)!;}93K -z-i84y>Pjp~)|Zl%Q$usfGz~R(Kg;Rj5-xKDa1UV01zzB2Vw__RA&o0c*?P)1@_%qs -zssiT3cW+v?EVM+#H>~CW!I|-Wg-EyJD?lD?*`op??S&&ujUm;sQt3up-1eUcC#WpS -z1}Yo+nVLYR(@a?L7hq84`O|Mg!tw3B1fFxsp#t2)HZjLflT%~j8y{}11}|=dkG)sg -z3$TRID&l3Prcl$FgiQq+$;3tRAdqrD?*c!(Qw=0S6nhBPuv4rz+2V`g(R=y`DAo{s -zz?idPOcJdkA0SfyOx=s75li`412Vf^^R0xbN!NX`D6C+wYw>b$97e?FU@22=^x972 -zmeU#XpVhLgrnkM$*EhK)SL)9~RO|_bgV`#*R}~t0D0)R44n-j$&WDbAc2XB_LtPT< -zY=#O1icY?94(cp7XOKBhB2!I7ZJK}_+o$kv-@gnVVI_ov-k$PS0MAigzYU{99B=*w -yS*L6N0oWJcoheb78g++S60=q_I7!el&_L+z*FDM95AoTZjGjeh|GhMOG# - -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir -index 5b06671a35..5f5a4c8a71 100644 ---- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir -+++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir -@@ -2383,10 +2383,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_17_0.mlir.bc -index 8f459cf76b62d1f070d0f279ab6943555b8338d2..8066e685e9032c25113740917eab9b7259813b80 100644 -GIT binary patch -delta 935 -zcmY+BT}TvB6vyY>Gq=05>$u*gRdOBU6OgfT_=68Y4F1PK-tkn`@83S7=sUE -zFf~qeXjG>m?S56ms#N{JJ1ZT}HV|bmjt@)dHi_@)nalFE$$$>#kifa{zYO)3zoTgGpxeZ4s>8&2t<>7~(jBBb-_mTJB)wH)z6ffVR{t!6J0YW_+nKA>;n4ItjkBN}&kXu~nQ}-A7KG7W1nQ -zZYnoFD4Dz7D!C#(%)$4pJ57GTnL*ocyI7sPoZ<%KEZdu8|BdnLRdW--YmhL{PN -z?wd8wh_-ZLgvc8&*W6n>U@3E}0FAP$n+a2s9&b;N*k0S|L5^WR@7K--k*B0M5;fSr*LjEZ9s3o_zHWT$VX-{Ny* -zo2=okHH!p#+@sMnskXfigY57XWU7g%og*NpHyQgGEWQ2-{(HWdt<#~~{aE)Bpy1wa -wXo67c-xSxYOtri*wu`xV3@rCY{1Es$Ius_`W4i~eJK$??UpiUt;`)W?UvT8H_y7O^ - -delta 935 -zcmY+BUr1Cz6vpqFbH~}c>#E~*{jULeP+qP_r$R=j%=8{EWgo*`{l8H$fNP&n% -zmWEmcjSoRDSyV4l!Y#^|$fq7eNHC#@gdTbfdI%wcnyV{rcVG^4&-b1A%{l8+us#I~ -zv&7vQ()AJ+3^Je3r;9nV22pVpyX5`>4;Mu%iL(9e#6VQksEZJ_zNuxZRQ++(Nfphg -zAQPpkMO({Hm^7JWa;sz}(zvK`)`|}28S<8h7tVvuxJkK8qZSvN&UZCZ&3VGqG$ZMk -zCQXw+G>w_8Ij0JHN)7_4t6( -z8xweoQ?nmU&g$x?aE7zSraO3-)5a!oA=gL41A`bA)4AOxaUFF|BN4lvG|zEB7njt3 -zbBkRjrx?T}XP6WfZWAKnsC%E&Fl>fl&cb5cUFA=5`sy_j2PjV5=-t< -zNQuwxm}(eki2S@`Ji*a}9*(0p%4rG^tiinRP<32!ry+?UX;%S}wn2)iA*3$c%Cule -z3>Td6$EZw70hNMNObsJ*(L|+^pN4kXjIRYgh*`f2E`j5Ul1jr2Hi*lGy=35`SSvie -zKim8`V?Mc}G^F7PnpuUPF*Syoi%HvA-A_iw#5Fx?hXwsdUi<1trK|&%G0Y@}Sn~L+ -z_eC#>wsoPG$Qy5#JS^`sl+_$SJ*(;g!qkY(+tnqSOWSOy+0uUK6Msra>$1_tcZAz) -zljJ|RX_<{@=;q4;Hr7hnxj;7mqn!e@MOjS4Qz8f0=JmjJao%e{Ox*S!aqRS4eDNNV -zHLP-9nn0J;=Z}$M!|pK14o@OeeTdo^0Xe<-z_0ww>!0K9^Tl$78``Z;6>k8N*1f7) -u@I-c}SgkVEu*WzdR)YaZSwDl_;Am}Cm~4+>bsG1) -> tensor<*xf32> { - func.return %0 : tensor<*xf32> - } - --// CHECK-LABEL: "type_per_tensor_quantization" --func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+// CHECK-LABEL: "type_quantization" -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK-LABEL: "type_per_axis_quantization" -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_18_0.mlir.bc -index 7ac3ab342767d1c4847bf8e4f17b6a584728e0e2..ed1b4ed419907bdc2e9db743333827bcece829cb 100644 -GIT binary patch -delta 991 -zcmY*XOK1~O6wQ5ip06|0*gmH*H6hfr9mi>mVT`ecSTsh1B9;1SwQ3bbP!TEOMu~zJ -zs#vKa4@FT>5XBFOI0$~I3l}1O7UFN$g$rE>f(th;#5*;`I0JK;ch0@%oVoLH1YV86 -zz4I(=ZE!3rD%#paJRVoq*eCE*1^emw0hcv-Nf!EhP1Y7qrL{-k`Cqk}8nNp19I<$n -z1=UKdHTwI8g3F6VL64+xk>}KWqk}zTYKt+T#*9Tq&gGI|)N=~W=d+0|Wx>^TyO^LR -zU8e(G7p_=3r4@J3o9=Gk&R9YHFdfT9S0dt|#6d*K6B4s3Ro_0R3%exJEqP}iuSh<8 -z0#8bw+Jz}8I`gM+REovLGk8|=#YI)AkMs1Ny_i;)>bFL77Fv>LGQm7*zAgbxNv<2i -z)6JCR*oGx3u%x`O3mLUJJlAmSz{3G4(n^I}7nCG_^aRHw0t6n&6W}4Ln2~5Bcz8}! -zzt7Oo;&1##96(!r3wxU5prd9^T_p<=BN&lb!HVRzjkbSys=->4uwi(IE7v6 -z#knXdddOIC4xq -zuriFw7wfyx|Mo*@MjyPwv_IbTj_DJUF-Dq?qgP>%+T0vx0|V+{b5CI4W%DjNn>VA< -zS8HR42`wE&`|IO4!&FV~5kj+p@T9M`?0;jYCkbQ;V>=D -z8?x$IY(?GV4%D>x5{l!mkCzz?_?HqMo9+aufl?nwp~VsPBMfMxpId&K)x>v^ezWnW -zHsd}qwgvwB{f6X2fRg{BwF4rVf2*CP3C#|6yHbhv7O;J%eJj*;cZb9j8tRYQ7eEc$ -dmRjiSHlJ?n<7OlrwL5M9j`KR=4|Tm4e*mxew6XvI - -delta 1006 -zcmY*XQD_ub6y5voJm1bt*7%&P8+TdaZZeLObzP=w8kS|7PPbhOQb}qPT?2~2V1w8# -zriIEXR3R24X$g;(77Hy%q*@ShC}^cW`6#8aAELH@e)yrFAoRl@Kg2t67uOkhmwD&h -zcg~snR_5W~dARZ&i|uwSDreYqIp|14Y)EV%N}iVJsa)5<)-skQ -z(kuDUF#aU@+*f!;a`zCrQVb2B#RVxUBj4bn`L+6%wf!_b6xuqMGGy-^ZD>R -zYF?6nrX<(vVjE0Xa_mQ63iOp1b0M!LW7`bd4n53Dky9#`ebblx`VIU;B0%VYJOLh~ -zig}4u1P{Lzt^afCOaD1OBMx9yJ&t*rZKJI=t=%C@65|+`*uz;{>c=F1Xioo9U~PNk*E#Cw_)$HQAeD^}70^cQiEc -zckdzEdt+mBqSMAHCbWDQ?WK>?4Bd|0^Mn>d;kK_2+3Uu>ZWG87CJ%zEK2D|+e#a4) -zRt4w?ep3CIw82x$$@in}8L0KCw<%8WQL4tE9-K>iY=a%921>$MRuSUCB6Pg{)c20?bKBxxHz&=p6V+G&a -z3lN=yh*(2vD(}pikM>M(GagIW!{y*t=Vsjabo%K2&!#__Z5&mL6)QMXSr#t) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_quantization" --func.func @type_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK: function_type = #vhlo.type_v1 !vhlo.token_v1>> -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir.bc -index 805452951125ee099c35ea3774a62fecaca96113..b8ebd6b5f8835df3a28981ba73e5c9d2828ce4b1 100644 -GIT binary patch -delta 940 -zcmZ9KT}TvB6vsJx_U>(VTID*n+s>@s#&Mm!jO)5C>$WXRXkcWJeLMsShLR+`ghYu+ -zf%sOC#En2y!lxjJKIW-J9(?erJG+o>44lKAbN_#S^FOn`0_!XA -z`~@}})G*FFMuIwyWm&HO79%uj+lDO==MCH6i1szWGvYk7H~5p^T{(! -zwxENN4xG@`j8XX@StL1uT26jQo(5ZT(`->!$N`%Ml~6&1X3LZ2h#D@Nwvom(L&R~0 -zx{Ug=rWow%-}4erScr!Vx&}-+_Lsp<`=NS -zs8|FPa>zQH$Ted#Bc2!7Dv$dNP-hDNnHASvMz9U5jAB)8SOPq`W0@7FhO(qgmX!%3 -zh8Acw4YdHV%E;;?Tw@3jkSS8YN!8I~$mK}tAaFUA%7+~X9l4r17Umqe|2?${P8wnq -z{zNdXBw;`#3|>--M913Xwd(Rq`v_!ZmZwx54yu2dUj?lQM=9LVu_(6=9mhNrq_8iX -z%M#_HoQQIY@1sa{;=(ADDBRZ3-7_@RP}8Z>!3-2BT%kBil`ardC2p_JOHh`Zb`J)f -zcKnW2U>+90Nx+1T75OyV(K=MWC}0r)>;#PIDEzWB8LAWZ2ZWE)OOTTTT_$=e(Xp-* -zHQpy(7uDOSd!vE+YXh~}Nd4_U>g`A)_0L`EPwqrv?+0(!eG%H=6@?ZZt9#Pesp}A` -zaw?NnLEp;^YX6oaKWB~v3HWirL*V)S-ZZp2K~kWSdZ_+rHT8l`; -z1T5FHc^6f`VA(k=t)5>{^)VL*!`oebQ?#~vF1vf%fcwL}r$BSS0UA->{#w2W8V6oG -Ka(}DulJyt8S(h>Z - -delta 954 -zcmZ9KPe>F|9LM?X&$q9cX_e>LZacF&jpH~z)^T08b=#ID*1*Uh!wx}$p(I(CkSH-J -z5N`!Zya+@kJOx4YXHL~65;~R;2pvj@{v3l29z6AC7t)P^_jvF3zF$7S?{C&uV0{Ii -zzrZG{>*pLjK^@05O-~QrVw9?uty>&%Ublh`wXadLd`>DA%8~dYmh4bX4nq-+mCIb% -zhAu{uhnp=<2cqpbp~zQ^a<~~050ZJ}b5wlsA$bNYVHrmFj2y6OxDv_plVxebh^~ZMML{LJFEO(E2-g?_gk+Kwa8mBb8J5Lx>L8TGR4UtG+h~i`)bWNg -zAOzo2o8UMQqwpt+j+8V2B4O~7QY0EyE3f4)zO)TPTKHBq-Yv}GN8nUQpRO(;`@)WL6oTXA1im4R0 -z)#`aDh)uf(!=1MOjul}Z7QjisgoZ`&)NgMYtUdH#5diE2jA`fvMQ1WnBkT_eAG=H7 -zi%O?~VWMN5C#$?qIxomu)w@wg{k4wTsHgt+AN6*$p8Dr5^(Sv4x3}T#nlC~tyrR&o -zVQEhqJ2f3bNlbOQGU$6<1IoYQh|gU|!vumjw<5rU{`54o*kMv2lX|F7L}e}la=LH2 -zf9yu|UadyNV;(HmvU#V6ncUl!fHtxyhzdQXF5gAFt) -Uy@R!E9+a{)3Xc$M^<6ao0<-j)6#xJL - -diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir -index 4f0822483f..49e87338c0 100644 ---- a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir -+++ b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.mlir -@@ -2388,10 +2388,10 @@ func.func @type_dynamism_unranked(%arg0: tensor<*xf32>) -> tensor<*xf32> { - } - - // CHECK-LABEL: "type_per_tensor_quantization" --func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor) -> tensor { -- // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1) -> !vhlo.tensor_v1 -- %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor) -> tensor -- func.return %0 : tensor -+func.func @type_per_tensor_quantization(%arg0: tensor>, %arg1: tensor>) -> tensor> { -+ // CHECK: "vhlo.add_v1"(%arg0, %arg1) : (!vhlo.tensor_v1>, !vhlo.tensor_v1>) -> !vhlo.tensor_v1> -+ %0 = "stablehlo.add"(%arg0, %arg1) : (tensor>, tensor>) -> tensor> -+ func.return %0 : tensor> - } - - // CHECK-LABEL: "type_per_axis_quantization" - -From a09e8f63d1a120d01fab2fd0bf879e8721f1d00e Mon Sep 17 00:00:00 2001 -From: Abhinav -Date: Fri, 23 Feb 2024 18:12:25 +0000 -Subject: [PATCH 5/5] Relax per-axis dimension check - ---- - stablehlo/dialect/Base.cpp | 13 ++------- - stablehlo/tests/ops_stablehlo.mlir | 45 ++---------------------------- - 2 files changed, 4 insertions(+), 54 deletions(-) - -diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp -index 299b1635f0..781d7c0a39 100644 ---- a/stablehlo/dialect/Base.cpp -+++ b/stablehlo/dialect/Base.cpp -@@ -67,15 +67,6 @@ LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { - return success(); - } - --// Q: Quantized (per-tensor or per-axis) --// NQ: Non Quantized --// tp1 tp2 Result --// NQ NQ tp1 == tp2 --// NQ Q false --// Q Q tp1.storage_type() == tp2.storage_type() || .. --// Q(per-axis) Q(per-axis) tp1.quantized_dimension() == ... --// Q(per-tensor) Q(per-axis) false -- - bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - // Get element type if shaped - tp1 = getElementTypeOrSelf(tp1); -@@ -101,11 +92,11 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { - auto qpatp2 = qtp2.dyn_cast(); - if(qpatp1 && qpatp2){ - // Both are also per-axis quantized -- return qpatp1.getQuantizedDimension() == qpatp2.getQuantizedDimension(); -+ // For now, don't match dimentions, Per OP verifier will do it. -+ return true; - } - // return true if both are per-tensor quantized - return !(qpatp1 || qpatp2); -- return true; - } - - return false; -diff --git a/stablehlo/tests/ops_stablehlo.mlir b/stablehlo/tests/ops_stablehlo.mlir -index 2d2ca1d014..3a0063e63e 100644 ---- a/stablehlo/tests/ops_stablehlo.mlir -+++ b/stablehlo/tests/ops_stablehlo.mlir -@@ -5393,16 +5393,10 @@ func.func @is_compatible_dynamism_dim_mismatch(%arg0: tensor<1x?xf32>) { - - // ----- - --// TODO(b/230263270): For stablehlo.add, the plan is to only allow fp+fp=fp, q+q=q and q+q=fp. - func.func @is_compatible_quant_mix_non_quant(%arg0: tensor<1xf32>, %arg1: tensor<1x!quant.uniform>) { - %0 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32> -- //%1 = "stablehlo.add"(%arg0, %arg0) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1x!quant.uniform> -- //%2 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -- //%3 = "stablehlo.add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -- //%4 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> -- //%5 = "stablehlo.add"(%arg1, %arg0) : (tensor<1x!quant.uniform>, tensor<1xf32>) -> tensor<1xf32> -- %6 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -- %7 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -+ %1 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> -+ %2 = "stablehlo.add"(%arg1, %arg1) : (tensor<1x!quant.uniform>, tensor<1x!quant.uniform>) -> tensor<1x!quant.uniform> - func.return - } - -@@ -5444,42 +5438,7 @@ func.func @is_compatible_quant_signedness_mismatch(%arg0: tensor<1x!quant.unifor - func.return - } - --// ----- -- --// The following is the not the exhaustive list of ops supporting quantized --// types. The list will be updated as part of adding verification support for --// quantized ops. --func.func @quantization_supported_ops(%arg0: tensor<1x2x2x!quant.uniform>, %arg1: tensor<1x2x2x!quant.uniform>, %arg2: tensor>) { -- %0 = "stablehlo.atan2"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %1 = "stablehlo.divide"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %2 = "stablehlo.power"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %3 = "stablehlo.remainder"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %4 = "stablehlo.subtract"(%arg0, %arg1) : (tensor<1x2x2x!quant.uniform>, tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- -- %5 = "stablehlo.abs"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %6 = "stablehlo.cbrt"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %7 = "stablehlo.cosine"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %8 = "stablehlo.exponential"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %9 = "stablehlo.exponential_minus_one"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %10 = "stablehlo.log"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %11 = "stablehlo.log_plus_one"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %12 = "stablehlo.logistic"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %13 = "stablehlo.negate"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %14 = "stablehlo.rsqrt"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %15 = "stablehlo.sign"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %16 = "stablehlo.sine"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %17 = "stablehlo.sqrt"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- %18 = "stablehlo.tanh"(%arg0) : (tensor<1x2x2x!quant.uniform>) -> tensor<1x2x2x!quant.uniform> -- func.return --} - --func.func @per_axis_quantized_ops(%arg0: tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>, %arg1: tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) { -- // %0 = "stablehlo.broadcast_in_dim"(%arg0) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x3x2x!quant.uniform:f32:3, {0.1:-30, 0.5:-20}>> -- %1 = "stablehlo.broadcast_in_dim"(%arg1) {broadcast_dimensions = array} : (tensor<1x2x2x!quant.uniform:f32:0, {0.1:-30}>>) -> tensor<2x2x2x!quant.uniform:f32:0, {0.1:-30, 0.1:-30}>> -- // %2 = stablehlo.reshape %arg0 : (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> -- // %3 = "stablehlo.transpose"(%arg0) {permutation = array}: (tensor<1x2x2x!quant.uniform:f32:2, {0.1:-30, 0.5:-20}>>) -> tensor<1x2x2x!quant.uniform:f32:1, {0.1:-30, 0.5:-20}>> -- func.return --} - - // ----- - From e4da62d2fcfacc4f81207163db8458b9bc6a3ec1 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Fri, 23 Feb 2024 19:23:48 +0000 Subject: [PATCH 03/14] correct comment message --- stablehlo/dialect/Base.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index ec0babf36eb..6caa6d45087 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -92,7 +92,7 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { auto qpatp2 = qtp2.dyn_cast(); if(qpatp1 && qpatp2){ // Both are also per-axis quantized - // For now, don't match dimentions, Per OP verifier will do it. + // Don't match dimensions, Per OP verifier will do it. return true; } // return true if both are per-tensor quantized From 7ece4541bc611686a41f23b1da7f3492bb56ff0b Mon Sep 17 00:00:00 2001 From: Abhinav Date: Fri, 23 Feb 2024 19:32:50 +0000 Subject: [PATCH 04/14] fix clang format error --- stablehlo/dialect/Base.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index 6caa6d45087..3732eba5263 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -84,13 +84,13 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { if (qtp1.getStorageType() != qtp2.getStorageType() || qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax() || - qtp1.getExpressedType() != qtp2.getExpressedType()){ + qtp1.getExpressedType() != qtp2.getExpressedType()) { return false; } auto qpatp1 = qtp1.dyn_cast(); auto qpatp2 = qtp2.dyn_cast(); - if(qpatp1 && qpatp2){ + if (qpatp1 && qpatp2) { // Both are also per-axis quantized // Don't match dimensions, Per OP verifier will do it. return true; From 52dee86277db380527c3037c134bd5c7fceca2b7 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Fri, 23 Feb 2024 19:47:39 +0000 Subject: [PATCH 05/14] remove getExpressedTypeOrSelf, unused after new changes --- stablehlo/dialect/Base.cpp | 7 ------- 1 file changed, 7 deletions(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index 3732eba5263..56c0da365d3 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -33,13 +33,6 @@ limitations under the License. namespace mlir { namespace hlo { -namespace { -Type getExpressedTypeOrSelf(Type type) { - auto quantType = type.dyn_cast(); - return quantType ? quantType.getExpressedType() : type; -} -} // namespace - LogicalResult verifyCompatibleShapeWithBounds(Type type1, Type type2) { if (failed(verifyCompatibleShape(type1, type2))) return failure(); From 290458a962e8c4ca4c4a67bb41e8edc0b960eb6d Mon Sep 17 00:00:00 2001 From: Abhinav Date: Sat, 24 Feb 2024 01:01:03 +0000 Subject: [PATCH 06/14] added comments, reorder returns --- stablehlo/dialect/Base.cpp | 27 ++++++++++++++++----------- 1 file changed, 16 insertions(+), 11 deletions(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index 56c0da365d3..83be8566f1d 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -67,13 +67,23 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { auto qtp1 = tp1.dyn_cast(); auto qtp2 = tp2.dyn_cast(); - - if (!qtp1 && !qtp2) { - // Both are non quantized + // Sparsity: In the most general case, we allow any combination of + // sparsity/denseness across any combination of operands/results, as well as + // differences in sparsity encodings for operands and results. + // Individual ops may introduce additional constraints. + // No additional code is needed to check this because of how sparsity is + // currently implemented. + + // Default case: Unless dynamism, quantization and/or sparsity are involved, + // the types are required to be exactly equal. + if (!qtp1 && !qtp2){ return tp1 == tp2; } + // For quantized types, + // a. both `tp1` and `tp2` should be quantized types + // b. with similar quantization granularity (i.e. both per-tensor or both per-axis) + // c. with equal storage_type, storage_type_min and storage_type_max, expressed_type if (qtp1 && qtp2) { - // Both are quantized if (qtp1.getStorageType() != qtp2.getStorageType() || qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || qtp1.getStorageTypeMax() != qtp2.getStorageTypeMax() || @@ -83,13 +93,8 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { auto qpatp1 = qtp1.dyn_cast(); auto qpatp2 = qtp2.dyn_cast(); - if (qpatp1 && qpatp2) { - // Both are also per-axis quantized - // Don't match dimensions, Per OP verifier will do it. - return true; - } - // return true if both are per-tensor quantized - return !(qpatp1 || qpatp2); + + return ((qpatp1 && qpatp2) || (!qpatp1 && !qpatp2)); } return false; From 810ec75ff0e257163ebfc10b6af65726619ffa1a Mon Sep 17 00:00:00 2001 From: Abhinav Date: Sat, 24 Feb 2024 01:32:32 +0000 Subject: [PATCH 07/14] fixed formatting!! --- stablehlo/dialect/Base.cpp | 12 +++++++----- 1 file changed, 7 insertions(+), 5 deletions(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index 83be8566f1d..d689f565ff2 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -76,13 +76,15 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { // Default case: Unless dynamism, quantization and/or sparsity are involved, // the types are required to be exactly equal. - if (!qtp1 && !qtp2){ + if (!qtp1 && !qtp2) { return tp1 == tp2; } - // For quantized types, + // For quantized types: // a. both `tp1` and `tp2` should be quantized types - // b. with similar quantization granularity (i.e. both per-tensor or both per-axis) - // c. with equal storage_type, storage_type_min and storage_type_max, expressed_type + // b. with similar quantization granularity (i.e. both per-tensor or both + // per-axis) + // c. with equal storage_type, storage_type_min and storage_type_max, + // expressed_type if (qtp1 && qtp2) { if (qtp1.getStorageType() != qtp2.getStorageType() || qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || @@ -93,7 +95,7 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { auto qpatp1 = qtp1.dyn_cast(); auto qpatp2 = qtp2.dyn_cast(); - + return ((qpatp1 && qpatp2) || (!qpatp1 && !qpatp2)); } From 13787aebe5a6b35e37cc9311efa69e0f2872032b Mon Sep 17 00:00:00 2001 From: Abhinav Date: Tue, 27 Feb 2024 17:54:24 +0000 Subject: [PATCH 08/14] moved defalt case to the end of API --- stablehlo/dialect/Base.cpp | 37 +++++++++++++++++++------------------ 1 file changed, 19 insertions(+), 18 deletions(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index d689f565ff2..5e43b8f1d51 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -65,26 +65,14 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { tp1 = getElementTypeOrSelf(tp1); tp2 = getElementTypeOrSelf(tp2); - auto qtp1 = tp1.dyn_cast(); - auto qtp2 = tp2.dyn_cast(); - // Sparsity: In the most general case, we allow any combination of - // sparsity/denseness across any combination of operands/results, as well as - // differences in sparsity encodings for operands and results. - // Individual ops may introduce additional constraints. - // No additional code is needed to check this because of how sparsity is - // currently implemented. - - // Default case: Unless dynamism, quantization and/or sparsity are involved, - // the types are required to be exactly equal. - if (!qtp1 && !qtp2) { - return tp1 == tp2; - } // For quantized types: // a. both `tp1` and `tp2` should be quantized types // b. with similar quantization granularity (i.e. both per-tensor or both // per-axis) - // c. with equal storage_type, storage_type_min and storage_type_max, + // c. with equal storage_type, storage_type_min, storage_type_max, and // expressed_type + auto qtp1 = tp1.dyn_cast(); + auto qtp2 = tp2.dyn_cast(); if (qtp1 && qtp2) { if (qtp1.getStorageType() != qtp2.getStorageType() || qtp1.getStorageTypeMin() != qtp2.getStorageTypeMin() || @@ -93,13 +81,26 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { return false; } - auto qpatp1 = qtp1.dyn_cast(); + auto qpatp1 = qtp1.dyn_cast(); auto qpatp2 = qtp2.dyn_cast(); + bool quantizationGranularityMatches = (qpatp1 && qpatp2) || (!qpatp1 && !qpatp2); - return ((qpatp1 && qpatp2) || (!qpatp1 && !qpatp2)); + return quantizationGranularityMatches; } - return false; + if (qtp1 || qtp2) + return false; + + // Sparsity: In the most general case, we allow any combination of + // sparsity/denseness across any combination of operands/results, as well as + // differences in sparsity encodings for operands and results. + // Individual ops may introduce additional constraints. + // No additional code is needed to check this because of how sparsity is + // currently implemented. + + // Default case: Unless dynamism, quantization and/or sparsity are involved, + // the types are required to be exactly equal. + return tp1 == tp2; } bool isCompatibleForHloTypeInference(Type tp1, Type tp2) { From 3209cd2bad6b4c2af3908ed591f3c36b2d2bd240 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Wed, 28 Feb 2024 00:17:40 +0000 Subject: [PATCH 09/14] addressed review comments --- stablehlo/dialect/Base.cpp | 1 + stablehlo/dialect/Base.h | 49 +++++++++++++++++++++++++++++++ stablehlo/dialect/Base.td | 3 ++ stablehlo/dialect/StablehloOps.td | 13 ++++---- 4 files changed, 58 insertions(+), 8 deletions(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index 5e43b8f1d51..05b7cc266fe 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -88,6 +88,7 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { return quantizationGranularityMatches; } + // return false if only one is of quantized type if (qtp1 || qtp2) return false; diff --git a/stablehlo/dialect/Base.h b/stablehlo/dialect/Base.h index 39174c7d159..ef8e3bb3b44 100644 --- a/stablehlo/dialect/Base.h +++ b/stablehlo/dialect/Base.h @@ -306,6 +306,55 @@ class CompatibleOperandsAndResultElementType } }; +template +class CompatibleOperandsElementType + : public mlir::OpTrait::TraitBase { + public: + static LogicalResult verifyTrait(Operation *op) { + Type expected; + if (op->getNumOperands() != 0) expected = op->getOperand(0).getType(); + if (!expected) return failure(); + + auto typeMatch = [&](Type actual) { + return isCompatibleElementTypeForHloTypeInference(actual, expected); + }; + auto allMatch = llvm::all_of(op->getOperandTypes(), typeMatch); + if (!allMatch) { + return op->emitOpError( + "requires compatible element types for all operands"); + } + + return success(allMatch); + } +}; + +/* +template +class CompatibleAllElementType + : public mlir::OpTrait::TraitBase { + public: + static LogicalResult verifyTrait(Operation *op) { + Type expected; + if (op->getNumResults() != 0) expected = op->getResult(0).getType(); + if (op->getNumOperands() != 0) expected = op->getArgument + if (!expected) return failure(); + expected = op->getT + auto typeMatch = [&](Type actual) { + return isCompatibleElementTypeForHloTypeInference(actual, expected); + }; + auto allMatch = llvm::all_of(op->getOperandTypes(), typeMatch) && + llvm::all_of(op->getResultTypes(), typeMatch); + if (!allMatch) { + return op->emitOpError( + "requires compatible element types for all operands and results"); + } + + return success(allMatch); + } +}; +*/ template class CompatibleOperandsAndResultType : public mlir::OpTrait::TraitBase; +def HLO_CompatibleOperandsElementType : + HLO_NativeOpTrait<"CompatibleOperandsElementType">; + def HLO_BoundedAttrInterface : AttrInterface<"BoundedAttrInterface"> { let cppNamespace = "::mlir::hlo"; diff --git a/stablehlo/dialect/StablehloOps.td b/stablehlo/dialect/StablehloOps.td index a1cf4f35466..ed3976ab8db 100644 --- a/stablehlo/dialect/StablehloOps.td +++ b/stablehlo/dialect/StablehloOps.td @@ -1532,7 +1532,7 @@ def StableHLO_TupleOp : StableHLO_Op<"tuple", [Pure, } def StableHLO_CompareOp: StableHLO_Op<"compare", [Pure, Elementwise, - SameOperandsElementType /*compare_c1*/, + HLO_CompatibleOperandsElementType /*compare_c1*/, SameOperandsAndResultShape /*compare_c2*/, InferTensorTypeWithReify /*compare_c1, compare_c2*/]> { let summary = "Compare operation"; @@ -1685,8 +1685,7 @@ def StableHLO_DynamicUpdateSliceOp: StableHLO_Op<"dynamic_update_slice", //===----------------------------------------------------------------------===// def StableHLO_BatchNormGradOp : StableHLO_Op<"batch_norm_grad", [Pure, - AllElementTypesMatch<["operand", "scale", "mean", "variance", "grad_output", - "grad_operand", "grad_scale", "grad_offset"] /*batch_norm_grad_c2*/>, + HLO_CompatibleOperandsAndResultElementType /*batch_norm_grad_c2*/, InferTensorType /*batch_norm_grad_c3, batch_norm_grad_c4*/]> { let summary = "BatchNormGrad operation"; let description = [{ @@ -1725,8 +1724,7 @@ def StableHLO_BatchNormGradOp : StableHLO_Op<"batch_norm_grad", [Pure, } def StableHLO_BatchNormInferenceOp : StableHLO_Op<"batch_norm_inference", - [Pure, AllElementTypesMatch<["operand", "scale", "offset", "mean", - "variance", "result"]> /*batch_norm_inference_c2*/, + [Pure, HLO_CompatibleOperandsAndResultElementType /*batch_norm_inference_c2*/, InferTensorType /*batch_norm_inference_c7*/]> { let summary = "BatchNormInference operation"; let description = [{ @@ -1759,8 +1757,7 @@ def StableHLO_BatchNormInferenceOp : StableHLO_Op<"batch_norm_inference", } def StableHLO_BatchNormTrainingOp : StableHLO_Op<"batch_norm_training", - [Pure, AllElementTypesMatch<["operand", "scale", "offset", "output", - "batch_mean", "batch_var"]> /*batch_norm_training_c2*/, + [Pure, HLO_CompatibleOperandsAndResultElementType /*batch_norm_training_c2*/, InferTensorType /*batch_norm_training_c5, batch_norm_training_c6, batch_norm_training_c7*/]> { let summary = "BatchNormTraining operation"; let description = [{ @@ -2814,7 +2811,7 @@ def StableHLO_TransposeOp: StableHLO_ShapedInterfaceOp<"transpose", } def StableHLO_TriangularSolveOp: StableHLO_Op<"triangular_solve", - [Pure, SameOperandsAndResultElementType, InferTensorType]> { + [Pure, HLO_CompatibleOperandsAndResultElementType, InferTensorType]> { let summary = "TriangularSolve operation"; let description = [{ Solves batches of systems of linear equations with lower or upper triangular From 4eba9128202f8bec57d3b3cdeb968615d4583524 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Wed, 28 Feb 2024 00:40:09 +0000 Subject: [PATCH 10/14] remove unwanted comments --- stablehlo/dialect/Base.h | 26 -------------------------- 1 file changed, 26 deletions(-) diff --git a/stablehlo/dialect/Base.h b/stablehlo/dialect/Base.h index ef8e3bb3b44..202b36df477 100644 --- a/stablehlo/dialect/Base.h +++ b/stablehlo/dialect/Base.h @@ -329,32 +329,6 @@ class CompatibleOperandsElementType } }; -/* -template -class CompatibleAllElementType - : public mlir::OpTrait::TraitBase { - public: - static LogicalResult verifyTrait(Operation *op) { - Type expected; - if (op->getNumResults() != 0) expected = op->getResult(0).getType(); - if (op->getNumOperands() != 0) expected = op->getArgument - if (!expected) return failure(); - expected = op->getT - auto typeMatch = [&](Type actual) { - return isCompatibleElementTypeForHloTypeInference(actual, expected); - }; - auto allMatch = llvm::all_of(op->getOperandTypes(), typeMatch) && - llvm::all_of(op->getResultTypes(), typeMatch); - if (!allMatch) { - return op->emitOpError( - "requires compatible element types for all operands and results"); - } - - return success(allMatch); - } -}; -*/ template class CompatibleOperandsAndResultType : public mlir::OpTrait::TraitBase Date: Thu, 29 Feb 2024 18:46:51 +0000 Subject: [PATCH 11/14] vhlo bytecode files --- .../stablehlo_legalize_to_vhlo.0_10_0.mlir.bc | Bin 20119 -> 20119 bytes .../stablehlo_legalize_to_vhlo.0_11_0.mlir.bc | Bin 20204 -> 20204 bytes .../stablehlo_legalize_to_vhlo.0_12_0.mlir.bc | Bin 20845 -> 20845 bytes .../stablehlo_legalize_to_vhlo.0_13_0.mlir.bc | Bin 20845 -> 20845 bytes .../stablehlo_legalize_to_vhlo.0_14_0.mlir.bc | Bin 20494 -> 20494 bytes .../stablehlo_legalize_to_vhlo.0_15_0.mlir.bc | Bin 16526 -> 16526 bytes .../stablehlo_legalize_to_vhlo.0_16_0.mlir.bc | Bin 16654 -> 16654 bytes .../stablehlo_legalize_to_vhlo.0_17_0.mlir.bc | Bin 17546 -> 17546 bytes .../stablehlo_legalize_to_vhlo.0_18_0.mlir.bc | Bin 17673 -> 17673 bytes .../stablehlo_legalize_to_vhlo.0_9_0.mlir.bc | Bin 19954 -> 19954 bytes 10 files changed, 0 insertions(+), 0 deletions(-) diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_10_0.mlir.bc index 1937d62ba1d1cf835f65d6c7f91007c2d3f7d02d..33cca348153a348af398f8a025096727aa7bac2b 100644 GIT binary patch delta 1019 zcmY*WO=uHA6!zuIblfcMI&8HhX}WB-Njk)xfBXTDp*lbDaD^E zUMeCoY7vPd;$2xh*n@bG-t0l4*Pc}HBG_9G)!B`-=`6c5JMVkn``-8F8#Ldbx9?%u z+_q17b{F@&WHOoaKSS`C>)NgaW^LDBi5=`P-GXGp!(4(D)$Vw2RMO3e8i^n4a3!y| zr7>>knBeMNyvB+WMp3QCM+&Yq*tWm_*Tx~BY1esdr^|yc0D{E@3hGayceH|vfPMrkmGMBZD_3z{u-q7Of|ttEm=U}>cmuZts}5pTz+_|! z{lX0BG}Af4L&1mc-oudK@)1F*$>i~UMGKZ-`FsG;%u0~XVEXonfCXEaHd$71UV9`A>Ur?gr|u(B$ARPTNG+3>p>|GdZOCX)lmI4ZfpahuVPsM7 z%5aWSL7j2BdkRcLG>mE3l7?}U#YvFGsm@XGm2;7d3LW6|)I>#@?je+@tWLUTM_V&{ zwuhNH8lx(wITN`S+O!TDCf9j1O%+a;IcB)77?kC@Y_=J%(x8f@`w{Hr^zZf|x=Z)S z>!K+WL#pV>c1e$j$nB!@rXG#ezt3-Yg-BCgJ&^Y;^Nfn>LN{^rKZtzE4*^Fb;4z^TVXt)=CN5w;AfsmFa8ZS_K5 z-Pjn6Vsd}L>DQ$bp}aqyzeJs0U~%P1$M2wrtlS$z+;No3v#%6GIrsF$A4lL=a`2T~@se42UQ=;x6kV z!|wWn7ez$MA4H;vcvl(^dJqrdO%E=3^`r}4g!MKDwK|b8R?}V8Uw!X=?|a|%F0OZR z{~^ezwV`qJ5O8H+WHr-~ZsBN|fhVR5l(>C|cNKxO0oxFG*zVMd zZKA<@az%fgIW_Z%OEw2=n+s!&p)h1>zS{We-`#9Hu1ByDSgTmh7eDN{houOiKJFxhQSLw5+scEB*UPOqon+g6ZBYeHfX> zS)4;OiUpU`rXJ~$7?%Nrqd4N)(RdSEl5SMX^yPI@dHvR5>`D3@&o>9V`qal0MLd!ici9?fIpg$NCR*mohO)K# z@rI7GJQySN)|4~p=pX}~%vXJ_88s1 z9WS$t6K^5_T1maW4!-_2`gPgf=-X!BrXt$OiA}rVga{#rEr0T&F$!W=J&m~Ys;{< z469Euu2D@pW@!m(S-P&<-Ye8;)HF4dBhF~1w-&8;s%C+c;bF>=xMF{@FU(;!B+e!e zb()-xM6x)pNR3hPIyp>Sj*2U~QgtvzBbDlspOCF4^(&#AaK#iiQ=?M+kTSI#<`^PQ zFq}YSB+0NKnBKRejAe%MVWe2XTZ~+}iq{yi8)!3HEH!b8(aO*?&M;CL5)ED7Qk|Q} zDx+Kga#hxmvgfKMh#0nf`59fpRuo!$n9c#k!3qqCe0Cr>k7vKfsxSxhU?t$X zidFHiknKL&YTzIMz)ZlTiribYPKR0~b)V!M3#7l}66-AH~q@q;uU`0bM zEtD3DhztrsOHt8-r@~&;g9k5Oyy(G$;Gc_lDT)e$Am~g2ZOX!X%)EWy{=WBiZ5h^< zVf86SsbQROj2Lws(=Lp`38d5;qegviKok8xA@Q5yu&h zBQg?am=(tQi6icFRnp?_q^H^ro z6M&vrFptLkC`K9K(nM3^4g=Jgr2k>|>MkRg!V05U5#31+TXomZivw6k){B zRIjF?rtT|@+<%CV7y>9VNeXyDX5>{Ep- zkt$M&NTq=;ic}{q)S*P-u8wY7afd8wI#qpL=@b^BBr@s2WpGUj zKR?}byw%`=0027%lRA3AqH{LXBB?2oKkmIgAENN46ENELu@@~8HO|$6!%f(~182Y& zr@hJ1cI0oHkaulH{<)3}4JcCRq&P*TRppmTzrR~`WyQ^mfh!g=$=*sEMp$clMyQDA z8CQ<;bLN1yz5lZnhiEkIW#-?mA?pDRSKyi J1m_F0=6|l0mU;jH diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_12_0.mlir.bc index 4260ea0e0ea674cd16f641895ae8b2daa2b39398..01bf1d1272e8931c8d1190c48c93d8cff04ac8c5 100644 GIT binary patch delta 948 zcmZ9KPfQa*6vp@Ubvo>>K!y#}c3ao&wk*?HYALmlwlO8~;sF90H8CD)G{iqw6eYx< zA($97I#EfC0WoUANwXI`=z)tDP99Cf#FGaxk%I>h=qzBUW|NuSefz%oec#))1z1~v zS4-GxP{TOx7-8x-rfIs~do-zL+lI{%rwrTsrVX~}c9xS`jdCP8aVU}ua=0E4=OV)` zHm3uj1fEu9jZv`@sgWEVu(1( za1@b|2*a#kW^!Kdave5Dl#G3+c z3Y2CF0g5Tv69h!UV42cj28oVa#SijEoV89sR=l$;xvu)Fw_(kKRs`iBg#{gpV%{FY zJQT#beWou@x4QmwOz2J=XD;zAvi5`~94x()TTtd~&Lso#7ju@_6wFJ32l{oVVT z(4eIZ(=Y>y6NalgmhUIKTQ~>+Q0y>V(vf?WfHl zyUxqE<6SpSHsXE%AFtsh?IXlN zc|#$sW3|D)ChG;?K0;Oe;#}sNOdr;E<|4+^1DjdAO`i%t#xt{zp~Lad26^@i*={1L z69JI3UFrVakk+c^rLIaMY$x1HEWP5y0%ROVDIpBaoL-G)>*82!F3#_eY%llHiW&Q!T+L}iI delta 947 zcmZ8fO=uHA6yBGY$#i#Z%QUnx*~E48)9oatF->C{k{}VhcxY)A6~seCDfn}&R!fm8 zrGlvFpcX`0D=PG)?4=&`(2JLzJPKCu!DESFh$a%2Ww}A`9a_||ouJJTr-HWk)fj9u?JOttI^{^R;$S!_E`*0WY)%gZ z;&@7zX+}9*^NE#kon$%cS9$9`7(Rb_p*tl4@Xz(OiB4CmAi4uHqCU#gb@P){gS@43-&n z1)wW#SRS`<*@{rlOT)) zb=9lusH=R1kta{_8AAZYOfm&LCr{)IYvOKnC)C7ZH1BsDbi}9VVSfz}ynV42pPK-! zm4c6ALiYFpkuX@M)Sp3O;-BJs`6ABP#~~x$+O~|VDta6CEEomQ_ET6eu^{H1A|IZ^w6I5lZ4!ykDK(m$(6K zWtfH;&=MiIVq*DTvbW6z0RWm4f{P})UN!l`w?$=8QCaI=gshl#W2pRFaz{3awf3Bq z+3}w1Cz|Pf`=4Ilzx39d=)nYNMG66mlT?aR3aM21m+cm!KH&qo z+m+s5O>V6k%B_^x$m!Ctxh=;&81gTyCad1?ULQckJCZL!C*R)wE`$1Awf&ly?{lE$ eE%%Lq;eZ1g4agy0z5gK?HHD_W!TVgiY5fLNrkSq* diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_13_0.mlir.bc index 8a673a32d22cf619e7487ce5d08e48ceec3351db..3648ae409c06e237212bc84dd220f8bb41c8da6e 100644 GIT binary patch delta 948 zcmZ9KPfQa*6vp@Ubvo>>K!y#}c3ao&wk*?HYALmlwlO8~;sF90H8CD)G{iqw6eYx< zA($97I#EfC0WoUANwXI`=z)tDP99Cf#FGaxk%I>h=qzBUW|NuSefz%oec#))1z1~v zS4-GxP{TOx7-8x-rfIs~do-zL+lI{%rwrTsrVX~}c9xS`jdCP8aVU}ua=0E4=OV)` zHm3uj1fEu9jZv`@sgWEVu(1( za1@b|2*a#kW^!Kdave5Dl#G3+c z3Y2CF0g5Tv69h!UV42cj28oVa#SijEoV89sR=l$;xvu)Fw_(kKRs`iBg#{gpV%{FY zJQT#beWou@x4QmwOz2J=XD;zAvi5`~94x()TTtd~&Lso#7ju@_6wFJ32l{oVVT z(4eIZ(=Y>y6NalgmhUIKTQ~>+Q0y>V(vf?WfHl zyUxqE<6SpSHsXE%AFtsh?IXlN zc|#$sW3|D)ChG;?K0;Oe;#}sNOdr;E<|4+^1DjdAO`i%t#xt{zp~Lad26^@i*={1L z69JI3UFrVakk+c^rLIaMY$x1HEWP5y0%ROVDIpBaoL-G)>*82!F3#_eY%llHiW&Q!T+L}iI delta 947 zcmZ8fO=uHA6yBGY$#i#Z%QUnx*~E48)9oatF->C{k{}VhcxY)A6~seCDfn}&R!fm8 zrGlvFpcX`0D=PG)?4=&`(2JLzJPKCu!DESFh$a%2Ww}A`9a_||ouJJTr-HWk)fj9u?JOttI^{^R;$S!_E`*0WY)%gZ z;&@7zX+}9*^NE#kon$%cS9$9`7(Rb_p*tl4@Xz(OiB4CmAi4uHqCU#gb@P){gS@43-&n z1)wW#SRS`<*@{rlOT)) zb=9lusH=R1kta{_8AAZYOfm&LCr{)IYvOKnC)C7ZH1BsDbi}9VVSfz}ynV42pPK-! zm4c6ALiYFpkuX@M)Sp3O;-BJs`6ABP#~~x$+O~|VDta6CEEomQ_ET6eu^{H1A|IZ^w6I5lZ4!ykDK(m$(6K zWtfH;&=MiIVq*DTvbW6z0RWm4f{P})UN!l`w?$=8QCaI=gshl#W2pRFaz{3awf3Bq z+3}w1Cz|Pf`=4Ilzx39d=)nYNMG66mlT?aR3aM21m+cm!KH&qo z+m+s5O>V6k%B_^x$m!Ctxh=;&81gTyCad1?ULQckJCZL!C*R)wE`$1Awf&ly?{lE$ eE%%Lq;eZ1g4agy0z5gK?HHD_W!TVgiY5fLNrkSq* diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_14_0.mlir.bc index cbb1219304759f39789a70f6db4ce8887836d0a7..04721ebc27531c085623435c54f4d41404da1733 100644 GIT binary patch delta 992 zcmY*W-Afcv6rVF^uh-d0*Xv}9yRO6RxUSc2+cpyw8T<&#q!d%@NkJ*pqK6<0l|EGZ zLNMY|K@n+BQh-D};!$3o|_86Ief=?L&s0+vo0iKaN zrWo1?S$s~l^)-35=@0&rHGsBQD9Hv~7hUnaq$^6-XS`@+jyO6S61(h=jpA zN&|8@Ca&c^nV1l5b~mKOtR0i_wTk{%dkTzUt%1U06RX9bb42eyC0;pw)oCJWBo&cV z<+(JFMrx9yZ?zf<_e@M~$yUm277de1WNW-sKYLCL#j|RA82`;i3iXJtxnUSHv7Smk zv0N1vjhWn{;}SQmaGh;BXigX|n9}@$Os&7O?2x{$s$Ho%Sw3Id8`D1K0EV{oeIxNH zrK+b#xXBJB{Csi{dd0Wo*f#kq+vV@uF8{|C`O_6^@)s%wASv>BEbMLR7c`U$QZE@5 z%{(DJ9OA9UhVx>PH_A)F&d3zdGT z^b5g=p@JgRqI&Qpy7b^fK|S@-L%jqQ(m$XN(Njh>qoL~zcV_OH`M&d=@661>%pAN} z!~nId02rJZ)u(?6BzyHU^L9rA{EGc!US_$QZ8@8AUb%HIHLXPF4 z%Mfvp;UFR-Wenq@BwW><#xz4Y8A+z_HX|b!ag>p0KSmi%rpECSqjmLHaDtJ#dYoct z5-^$k&{TAiX&CShqm9>YU?Zb30T>gv!bckt0Sqw0WxZ?6cNs_ok&`|LVBlA!2DMJ8NfGh~`wA_(1 z%p#=lIn~z3WA8#Ko)~k?U)M7haF`9Q@i8 zr-Pz7Sx|aR;-(d@y-f$r3BmCp=JQ0qKHR=nY?^{L%AUJk`d9u zqte46-fC<(FBW;TEZSS;`*rB@ns!aXPDeQi(!pmE8%8w7b&$RLocOuv%aVVqFJg7K zvAun4-I6bUjCsY$T7Zm~*;5Ope6zt7e(H;=29Judlnq($L8=c5aygAkqnwxN%TO>1 MYd_#!-*+?o4F|7@zNb^L#tb%J_C@cWv60aa@LV+m=NZHCb#47YockvDP6Y3o{P|1{P{o zTV!Hj&_kjkbP1_rFxNUnhYnqeP8~aRwIGD>5WU&W+GgO*pYQvA|K3|I!)h5`UL!&C zgkwgS<5-sEdduWf17yT0%_U1jl(jSBtJX&|Nz^hT*|=pUFu@UVisKX_rxwSw=+N8u z=P<`n-kfGT@jj=sS8oo)g$J0E7*JIF7|Sk6o5%l(MP*8 z1{$2EQa-H2SmS^?r}FPZjs}`sH^DIc7m7YA#a7I7hIx?+CJ&}Sj%ns@10ZlZ4fsTY zI90({8a7P*fw8h(CR&^Y7#LwVy z)wa?0hMK-W+YzOy2n9^Y9#zuQVK~fG7pacbN>po=+@PqF*Spl&L$?em0S_laFNL& zf0D_`%|F$PGolbWK;&a4wADNtB72GAjKMqX^}EG}6A?>hvp#Z^fMSoqdrW#2^8(ST z${LB377A~~$79OJ20){ZzJo9|?Q1KQ#E0;(uf8oZ1><5RGMB6OaCx_}X5tF@&&9l~ zW@|P?cQoz`JsCX*u6PxVMe=ogkE<}!FP#(LqY0T~xMd))8@)K+a!Ag4FItKOO5Qi? zHiR8N3y@iyMW$Ma+5(Y>KrF|$`QIE?ur7$rI1i4@*MfL}{%JJ3#ZLSOOOAa~oI*vNR!vgb9WC|%*l=)C#VR4$2 zMJfgc4H6ZhM@T&e`&vQt&_gdpPd!HO7K9K!M6+JjwF|S$eBbx`H?uMiEA#N;ItiR| z%rJ8t%d(QDw?sbGL0X*Fl44OCBMU@))rvIjqLvnJ!=~w=!x3?U;{+n77RQum(_8m< zULK<>^W<~U0Vz@(_?gI#F@ z4NhIxhm{zs98l*}?mpmXpviR;48#AS=%Z3>!8~V}7pZ`I&;>cBnY#^u!09C5V+mp> ze8Ep7)taQ@)YNxFco5f$aEe!z{#M?kdwQaP;&%g<_ zAhEr^#xD@>RvJrCL`U|hlAiX%5vC@Q>bO(sMn?=9#~~#uh9%cm{_1@;ECrbzX9+5psHr_=YCCE=8y(+LvL#r>IVOkv zNhT*Z{!}L~iem5pk)NsHX46!V>?Mjb1n)5CcZ+o=EEdcrz4sUa#U6t9nDi><1)}XJ zYamiuFti>Ui7FrK0F4^@Hp0}juWevJdf_-F7!k|ia<9qJBNyr>S%{}$I(TnrVhvlsIytzbR!24$1 zhLG)N0Wym-$W#kan*pB^xytmoX27BCp o2fQsI6aER8McQtGg4b)0fylI)-s|`ks4s)UWM1IqT2Jo#4NEziUjP6A diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_16_0.mlir.bc index cbc5d43c644a58e1da5049e2ec5aecd5b89150d5..d316a36453f0aff77fbcbd33193b290ba0caf32c 100644 GIT binary patch delta 876 zcmYk4T}V_x6vy|>x#M`fZrj=Gy4rQ!uGeMRwslbv6IrSxJ_s6R4+UzGTgeVT=%QkZ=%@xD~*ZJZRMR zG+`5A00YG98}SbD(W^K{JU)PN67`MOaDqg0>l9um-rS0_ga!wL)Q5&wPE{HLyi20( z#w~0kk&u81`OY}jk|reL$)uN}D&Hee3a)k^r~sN|m_Xz$IC{M}MlnMIGqSTZF_i=* z)6J~}Kq&Niq>4BN8;@11r>YfkQ};UPI_e%zhWPU(TqXqYTv=*?!-N(l32lUWToJy_ zF*BX}gBuD3&?cVEVm3rPm%|+KcKNf^)*J^N8S?jNwvD!X!M_SMhkX7#`btVT^x}U!8-5vD!8_LG|2Wa zpVJ4QuuUisnKP}hP&l83cH2{RA7;FMf1w^(nl z+!gNf*!shRYVTusEV`3$oCz)E`P)RCWEhwI6=j~_N=3DYbHCyo#N|dsB%Ikn|DvE% zAbd>LRmLDCk5_gUZ=J7vS$RNd-POuGgI+gPH4EU#H<2UFv6HIpEJ)R^u} H?R?c=ji#IJ delta 892 zcmYk4T}TvB6vy}6GuP`lZrj;)U2WX9ab1>aTSY}w7O_d(EHE_7>Y>siM9e_QWDP%3 zD#9@GQbG}ZnD`V_^B(mQK}1hI_hxUB54{*fch*b?X6DX0=YRgcbMMU0!2Arno@T|T z;-QF$$4%2rhTOYsy#Ug3Qm+)f%tmL}mVHHG`CdQ7Zn1?Tho^1S(=uDM7K&q>5OIKT z0Fk&Ez?3{}RQ1$jJz)R?#A|Bt9`VsJ%n?uYV}e9Y?RA_W(a<=BH;6Yh;w+)T!60>` z;k8qp1_AGrXuf$Hn@J=kU{by_jyI(ViMVZhC93lS0;S+8x~LKuA|nJu-h!j&#W98% z5}1)4#rBj9N~VXl3IHMR^H?>p5}xl9W%X295fAB}gRZ0Q#WTd8FX1vFfH#$?7C1y` zqD|O>P>UJAZIPp#ZiJZ*9faFmcC0hj^R(S!`*Rg_aEad$bm1?_Tz=LS?(p zp9KfwN~0-4-UV$!OCmQgpJ~Uqyc#$bO>vb}1zZ*M2rY%e;0;kUuzcTuwIVP1$MSyexuWMJH<=?#792(8ndvBi5-nsgjR#PA(ZTVu!0Q^%ZXsM z_j(mv+dUFwdzdfY3!kxBC=ls{%!p7pUxaqYQ*|C?;In$+D<&6WO>VfPj_G-tAX7M9 zrO)7uTrV|QcemUX>GIh6A_FS+DUyrrBphc#OMCNeB2F?)$iA`?PjIbluZQ!n>>?!O zMp-nH*+KszuTvm=R92TKASF+fcNA{zuY6g4P-)%O@+^aHH(fCcV97Vp67~ET{YC%B zTfpu4yWV;8qh6oW=nRuO<8ypY(=^FsI*Gj|ha?oGKMWL>?NQ!D@8Isl~r9_dapU#v-(L4(B zUZR@R4@GH%nkJK`iWx{#BIK%bza$1s5d+OJLv*<{(dRP7Bi9jE%Aj0gQB8^Ot~a$Z zmGiWrs=67H45}(0RFxU*SV0*3!@4Zg~L_=Z~K%||JVM-LK8*`CXG(=x<2h@wD;-Iu^wza+%UjpZ-Ba?-BY!jEtf74HCHr1WYrP>y}g6FW1|32+rivKm=d?k_VkDYWu3OL zsj>kW6yM6mQaj~MZ3DOT&5&IJ=VY`#uW+Hf!j|@``~t*8Wgz6w?5Ou7MQL9zI4$ii=+)cMu IYgR)4062THrvLx| delta 939 zcmY*WT}TvB6y9^s-0sf)sMm2@-DO*K9oKQ&wrw$67qe+g0!fUr*djqmM8X8c2Ma+B z4V7NHPx+8VMbJYCLzG1c^$?^!DD_%z5fMQTy+lNJbft9$&Y8LAJKy>4cV=f`b_N!2 z5%0OErukVU!h*q|CKkvBB*YNrWL(8#Wr-@{*m|~eI3a43B}iDGlu|{aemWZ@isn&} z&l1(7eke{G)HIn4Rm?z|5+PT;Xmed4FNk>HI_ydrl*=-zDe>L)x=x}wPaCSLn<3ex zs`5ZpnZb_ks>UAqryAYegrvoWyH1N@lq2FG$3a9+LmZRB>#5nB!3;{w>h1- zj#Hc(eQ0pj(0Bu9IBRa1!<(Eow}^|LAdU53#+VrMbR9}*sBs#P+kCQlfdjfZm8bVO zYUtx`A84BWLD8n87(tUWG==6hATFA{`&_ziiyG!ECPuu~2Te|`r}&H`fGtXj1RRwa z;~cXHNnB!i!zRbgzxZ2X0JCDvn^bfibwP^SJZX+TduA38tX$C-s6OfNwt|UKiB$lR zc0iUXQKW9%tQ^Ov=q_o8q*yEo%D8Gr>ubpsaE>^#tuTif(N)?*`Uk|D((?zBRB}ij zDtYcRC5FsRPpvsJt?&?yg4!ddBvExU+19Rw$mp0D&=PjKqx~rAT>t2(I14KnW0Fa9 z`26BenbTS-TO{h{&;XHBhx~W<^y-fFU4U8va}!}o+#Z|DiT&jrwy??aKIj$S%C9vP z%A4Ef`OSiPk2h)LRp^8`6VRnaR|6-U+pdcD zfg`fl@&sE6+^)0hTG!EZT{d-fn#P}&W!cO&scmy9qEf=bKtU^b$psbw+n-n+x~Mz2A4v`R;cnr(kjl=5Dco zdCWG=GEq|_LZOhl&z2yj&S59*i@3KWR?W1vFBdCg%F*W`=6=?LIP6p8IzC5)PLHqmfWQd~df9Tgs$k7*;AwNrpie zh9MlWXI%t)>0~%PJ&YyQ7r$*f=tx8yk~oAYd01jxZ7Qs7Xu%eVv?On8#=DYV8ON)V z*Y{z)6iv;OI4#Al)OEZm`L2{2EDZ6+{?pi~ZWMNJPnc**9*uhSDEXlT^h-{+GZIZK zmHtvNP49=JSB_&1rlr8Nstq_0RfhsweYWL^Ix9t^ng~R;r6qqkkFO;Hc%n!o;0Q&G zO0*DSs6^JVu}i5H{7D=@OMME&blXN-nMDT*lM+k4o|c7{d#cFCBJElr11WTfs{U;6=_Wb0B+2oG*S~3GOSf5vTR_HM_{9<4@!nEcL+`(xi`aaIWqnFhK zGs-A@FqeIui%VMYD7+)z7piL9<&rsOtcs{$q++)0-ll#_TU-Zd=5^Y^gx=`Ap{q-s zEN}CKzAW#9epMN~kjR%imLsm|onZe7o+329rdTZmD?LeHgC`RCZ+eg`z+1KX3_N0( z&eQQ^6{RQhaWz?CL0mnn*jJF7iTYEqn`*kLP=-O5do3Je8*R@GAU7sZ=pjU7lmQKQ zq4JwQ>%>!BwW0w!-CtF&0MhP@NF5}i|0cXy7rNz*c)!+8#&|0``??$g+F?R0x1_ck X3JxCB1f6X6ymb-O8OwD2_Dc9KH!HIW delta 1026 zcmY*WT}TvB6rOX>T(7h1TDQ@1T~>3~aU0id8J4X!DVsKzASxv+wuhid!a{wJhzOw! zEtTrwM$ki0iAg<_G(<_zQ*RL>) zbR=P#8$>)VB9Vxx?y+S^t8-YSeF^u}rCXS``jumITG{#nq`i-NSSQx6>MXH%i3O!d ztb*}=P0r@7D`;yGHu8+J1MOmh+31ei$_du1kw6!F#ngOYXCPy9NmMj4>U-ezCd!mD zXB&p)MyY^d(1l?LTkNii;R*UP?BmB7%c)PngqgsEM8sicNi`9QZE+=kzJM@&fZI4NsHNXm zP~V3o`2484Z>AZAcjo6n-_o)cJ`8Wj_r;BI?{nQeGgd@2GEy0qVu`+R0?*#i#@D!n$RmEyC-0VyG5LyH9@l!w7=M*!-(E&Q dJ4i_0t7|<5>U*2%jaws7Jr0`i6P_lFroY0!v^xL* diff --git a/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir.bc b/stablehlo/tests/vhlo/stablehlo_legalize_to_vhlo.0_9_0.mlir.bc index 6e3e0284d5ab60a7e9367ebc5b68e05ac184a6c7..e7fd169038e0d9780b5c55da7f23bccc0faf1d3c 100644 GIT binary patch delta 1005 zcmZ8f!AlfT81M7v+vn`GF0WD7-E|wsab2I=x^A+l$r^S@0#VW^$U|s&46#rtDYZi& zWfAWXCJ|jabdYE)DymEU4IVmp2rBAeV000R-t0nZGtBqRyzl${-tYas*;=8k6?(dg zu*EIwjB7=>>&D~pl=l|D13|~J90@E~j`tzhSZO*L$$ERa1k0+QwT}LwIV(+q&r!)!~Mv(WYTa{Ha19(nu|GB5=jAI zRA3YkEGm#uUaWCAkGz1s1C=Si-a@9QI!r{Jj_w zC3U8Co>Ef{?-M-RvzG?0Fi$HeIUZP&PWT`pd23Y zeWmOAeD5D@>mGuw`s;G#t_xQ!)b*6xHf-;G-8#8R3UT@#KvFlA6Tk#L;T+6Fv9hYO zdUJ4)GHSnFr=#?vcif&JBTq*;T{Dqa@9O)Iqjpu{oE^?E&7zr@W~!{<;3iAa2z79} zXd<<%ct95olWQDJQJ&Kn#|+mvAIvqlrd02sF7@1Lg`dATKSJF!M+@Xe=(LG$bt}_Q z)l+Jev_wQsgoaE=uhtzeD=BR872ZfLQ&!bC#gNt}7n)8K^)5D@)2rW_ruV2`-=n&6 zuj-Hgsoo6iRlTvJ`dMnMefNPkOT3^2z2X!wQP}0iRtX0xsJ`Z;Hu^&IG2`EG)XU~0 zz60K`X1k2iUUhnss$AbG(oWqrIAF|{k)H1TL!WmndL1lUBtw#xOWwQ)>ff~1?Nr;D xuT=0pwbqc8-Wh!$SRB2r%(evCo@^T+!zGt^Ky!QZxda(D*~L%WdwjS*{ud1Fojw2n delta 1014 zcmZ8f&ubGw81>8RWZbN2n6BC;o49VbX)?w&jUmJsqSS*$5R0uS=s~RXD59dGqHR0~ z;xA+lQj3UJ52CJ$ih9w%!Gi}6f{Gp#DqaMC%qC(RXW5zI`{uoQ-}`26mgZ*Z?j0n! zW$p5;Uha9eZRf1uIetcCj%zs*Xj)G2DmKtRaD{vXk3;9LYCD`!QkY zu;7x8Mbyb;h4~VOuO1|~`HnQ$qTheVzfH&z*I2g9<>Ab@WYTa{F*Qn#no7Br2TuUl zC$J9?EGbY>L26*UjIx0K1S^(sUa*6Qa73`|9%KbCmX6}M;FanLoD!^3#c2VPkty^K zGo;f@`w15WuT3>k6Fhf`kW;Jdv2}g|3BjbUx5a!#kVY^ay(*A^B}|JFiEzUZ-o=nX zU2xRZ16z_*Z|&3$Hyd7YO|#BuR>3SIBr^;cVZ1Kb>XUJdP7%(K*7uOYw|pbbNU>Grw?@zzy#gq9Lz*3b4Rc0 z*@{srs3lHX&oZ2Y&CWqG%CwQw853poV&w>mR8ouG-QxwOSu_*VOx=q)xXH4#hpLF6i?nBW;DNLWSqDGvuq4gDHe2HwL$~^qwEwtD~O=k9DYi-J!a> zQ}vtwRDX$fs(w4K`hISrys+ReZC+4@o^Xnqs4sA1u8jlL)kr?8jXsdyWc+K6dYoS$ zIuLx%I}wt>vce%+>V{5{cIvXh0pnDJba$_>em`%~(^$(Q8PCy7+nYx*{Y?A3b?S%z wk?O&_HOtA8^R17@TCHy=b0|ZNKn`sq!zGt^RC5Q@r3_I+-SVv9-r8aN581Y@8vp Date: Thu, 29 Feb 2024 18:55:50 +0000 Subject: [PATCH 12/14] fix formatting!! --- stablehlo/dialect/Base.cpp | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/stablehlo/dialect/Base.cpp b/stablehlo/dialect/Base.cpp index 05b7cc266fe..7cb3e865763 100644 --- a/stablehlo/dialect/Base.cpp +++ b/stablehlo/dialect/Base.cpp @@ -81,16 +81,16 @@ bool isCompatibleElementTypeForHloTypeInference(Type tp1, Type tp2) { return false; } - auto qpatp1 = qtp1.dyn_cast(); + auto qpatp1 = qtp1.dyn_cast(); auto qpatp2 = qtp2.dyn_cast(); - bool quantizationGranularityMatches = (qpatp1 && qpatp2) || (!qpatp1 && !qpatp2); + bool quantizationGranularityMatches = + (qpatp1 && qpatp2) || (!qpatp1 && !qpatp2); return quantizationGranularityMatches; } // return false if only one is of quantized type - if (qtp1 || qtp2) - return false; + if (qtp1 || qtp2) return false; // Sparsity: In the most general case, we allow any combination of // sparsity/denseness across any combination of operands/results, as well as From 2b017303651b048b3c663dd288d5c101a4a2f949 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Thu, 29 Feb 2024 20:59:01 +0000 Subject: [PATCH 13/14] addressed review commit --- stablehlo/dialect/Base.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/stablehlo/dialect/Base.h b/stablehlo/dialect/Base.h index 202b36df477..23e1103b944 100644 --- a/stablehlo/dialect/Base.h +++ b/stablehlo/dialect/Base.h @@ -325,7 +325,7 @@ class CompatibleOperandsElementType "requires compatible element types for all operands"); } - return success(allMatch); + return success(); } }; From 2a8fcd4b0db6ce1c6ef579a674ac85e3479ef400 Mon Sep 17 00:00:00 2001 From: Abhinav Date: Thu, 29 Feb 2024 23:38:07 +0000 Subject: [PATCH 14/14] use verifyAtLeastNOperands trait --- stablehlo/dialect/Base.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/stablehlo/dialect/Base.h b/stablehlo/dialect/Base.h index 23e1103b944..436ee587f7a 100644 --- a/stablehlo/dialect/Base.h +++ b/stablehlo/dialect/Base.h @@ -312,10 +312,10 @@ class CompatibleOperandsElementType CompatibleOperandsElementType> { public: static LogicalResult verifyTrait(Operation *op) { - Type expected; - if (op->getNumOperands() != 0) expected = op->getOperand(0).getType(); - if (!expected) return failure(); + if (failed(mlir::OpTrait::impl::verifyAtLeastNOperands(op, 1))) + return failure(); + Type expected = op->getOperand(0).getType(); auto typeMatch = [&](Type actual) { return isCompatibleElementTypeForHloTypeInference(actual, expected); };