|
| 1 | +diff --ruN a/stablehlo/BUILD.bazel b/stablehlo/BUILD.bazel |
| 2 | +--- stablehlo/BUILD.bazel |
| 3 | ++++ stablehlo/BUILD.bazel |
| 4 | +@@ -1577,8 +1577,9 @@ |
| 5 | + srcs = [ |
| 6 | + "stablehlo/conversions/tosa/transforms/StablehloLegalizeToTosa.cpp", |
| 7 | + "stablehlo/conversions/tosa/transforms/StablehloPrepareForTosa.cpp", |
| 8 | +- "stablehlo/conversions/tosa/transforms/StablehloQuantLegalizeToTosaRescale.cpp", |
| 9 | +- "stablehlo/conversions/tosa/transforms/TosaRescaleLegalizeToStablehlo.cpp", |
| 10 | ++ # TODO: un-comment the following once #2751 is fixed |
| 11 | ++ # "stablehlo/conversions/tosa/transforms/StablehloQuantLegalizeToTosaRescale.cpp", |
| 12 | ++ # "stablehlo/conversions/tosa/transforms/TosaRescaleLegalizeToStablehlo.cpp", |
| 13 | + ], |
| 14 | + hdrs = [ |
| 15 | + "stablehlo/conversions/tosa/transforms/Passes.h", |
| 16 | +diff --ruN a/stablehlo/docs/generated/stablehlo_tosa_passes.md b/stablehlo/docs/generated/stablehlo_tosa_passes.md |
| 17 | +--- stablehlo/docs/generated/stablehlo_tosa_passes.md |
| 18 | ++++ stablehlo/docs/generated/stablehlo_tosa_passes.md |
| 19 | +@@ -10,17 +10,3 @@ |
| 20 | + |
| 21 | + This pass adds rewriters to make StableHLO ops more compatible with TOSA ops. |
| 22 | + Currently simplifies stablehlo.dot_general into stablehlo.dot for easier lowering. |
| 23 | +- |
| 24 | +-### `-stablehlo-quant-legalize-to-tosa-rescale` |
| 25 | +- |
| 26 | +-_Legalize StableHLO Quantized operations to TOSA rescale operations_ |
| 27 | +- |
| 28 | +-This pass rewrites StableHLO quantized operations to integer operations |
| 29 | +-by inserting TOSA rescale operations at the inputs and outputs of the |
| 30 | +-integer operations. |
| 31 | +- |
| 32 | +-### `-tosa-rescale-legalize-to-stablehlo` |
| 33 | +- |
| 34 | +-_Legalize TOSA rescales to StableHlo primitive math operations_ |
| 35 | +- |
| 36 | +-This pass rewrites TOSA rescale operations to StableHLO primitive math operations. |
| 37 | +diff --ruN a/stablehlo/stablehlo/conversions/tosa/tests/BUILD.bazel b/stablehlo/stablehlo/conversions/tosa/tests/BUILD.bazel |
| 38 | +--- stablehlo/stablehlo/conversions/tosa/tests/BUILD.bazel |
| 39 | ++++ stablehlo/stablehlo/conversions/tosa/tests/BUILD.bazel |
| 40 | +@@ -55,7 +55,18 @@ |
| 41 | + tags = ["stablehlo_tosa_tests"], |
| 42 | + deps = ["@rules_python//python/runfiles"], |
| 43 | + ) |
| 44 | +- for src in glob(["**/*.mlir"]) |
| 45 | ++ # TODO: remove the following excludes once #2751 is fixed. |
| 46 | ++ for src in glob( |
| 47 | ++ ["**/*.mlir"], |
| 48 | ++ exclude = [ |
| 49 | ++ "legalize_quant_ops_to_tosa_rescale.mlir", |
| 50 | ++ "legalize_tosa_rescale_to_stablehlo.mlir", |
| 51 | ++ "rescale_interpreter.mlir", |
| 52 | ++ "binary.mlir", |
| 53 | ++ "nullary.mlir", |
| 54 | ++ "unary.mlir", |
| 55 | ++ ], |
| 56 | ++ ) |
| 57 | + ] |
| 58 | + |
| 59 | + test_suite( |
| 60 | +diff --ruN a/stablehlo/stablehlo/conversions/tosa/tests/lit.cfg.py b/stablehlo/stablehlo/conversions/tosa/tests/lit.cfg.py |
| 61 | +--- stablehlo/stablehlo/conversions/tosa/tests/lit.cfg.py |
| 62 | ++++ stablehlo/stablehlo/conversions/tosa/tests/lit.cfg.py |
| 63 | +@@ -25,6 +25,8 @@ |
| 64 | + config.test_format = lit.formats.ShTest(not llvm_config.use_lit_shell) |
| 65 | + config.suffixes = ['.mlir'] |
| 66 | + config.test_source_root = os.path.dirname(__file__) |
| 67 | ++# TODO: remove the following once #2751 is fixed. |
| 68 | ++config.excludes = ['legalize_quant_ops_to_tosa_rescale.mlir', 'legalize_tosa_rescale_to_stablehlo.mlir', 'rescale_interpreter.mlir', 'binary.mlir', 'nullary.mlir', 'unary.mlir'] |
| 69 | + |
| 70 | + # Disallow reusing variables across CHECK-LABEL matches. |
| 71 | + # A variable can eschew this (be made "global") by prefixing its name with $. |
| 72 | +diff --ruN a/stablehlo/stablehlo/conversions/tosa/transforms/CMakeLists.txt b/stablehlo/stablehlo/conversions/tosa/transforms/CMakeLists.txt |
| 73 | +--- stablehlo/stablehlo/conversions/tosa/transforms/CMakeLists.txt |
| 74 | ++++ stablehlo/stablehlo/conversions/tosa/transforms/CMakeLists.txt |
| 75 | +@@ -25,9 +25,11 @@ |
| 76 | + add_mlir_library(StablehloTOSATransforms |
| 77 | + StablehloLegalizeToTosa.cpp |
| 78 | + StablehloPrepareForTosa.cpp |
| 79 | +- StablehloQuantLegalizeToTosaRescale.cpp |
| 80 | +- TosaRescaleLegalizeToStablehlo.cpp |
| 81 | ++ # TODO: un-comment the following once #2751 is fixed. |
| 82 | ++ # StablehloQuantLegalizeToTosaRescale.cpp |
| 83 | ++ # TosaRescaleLegalizeToStablehlo.cpp |
| 84 | + |
| 85 | ++ PARTIAL_SOURCES_INTENDED |
| 86 | + DEPENDS |
| 87 | + StablehloTOSATransformsPassIncGen |
| 88 | + StablehloTOSAPDLLPatternsIncGen |
| 89 | +diff --ruN a/stablehlo/stablehlo/conversions/tosa/transforms/Passes.td b/stablehlo/stablehlo/conversions/tosa/transforms/Passes.td |
| 90 | +--- stablehlo/stablehlo/conversions/tosa/transforms/Passes.td |
| 91 | ++++ stablehlo/stablehlo/conversions/tosa/transforms/Passes.td |
| 92 | +@@ -33,24 +33,25 @@ |
| 93 | + let dependentDialects = ["::mlir::tosa::TosaDialect"]; |
| 94 | + } |
| 95 | + |
| 96 | +-def StablehloQuantLegalizeToTosaRescalePass : Pass<"stablehlo-quant-legalize-to-tosa-rescale", "mlir::func::FuncOp"> { |
| 97 | +- let summary = "Legalize StableHLO Quantized operations to TOSA rescale operations"; |
| 98 | +- let description = [{ |
| 99 | +- This pass rewrites StableHLO quantized operations to integer operations |
| 100 | +- by inserting TOSA rescale operations at the inputs and outputs of the |
| 101 | +- integer operations. |
| 102 | +- }]; |
| 103 | +- let dependentDialects = [ |
| 104 | +- "::mlir::tosa::TosaDialect", |
| 105 | +- ]; |
| 106 | +-} |
| 107 | +- |
| 108 | +-def TosaRescaleLegalizeToStablehloPass : Pass<"tosa-rescale-legalize-to-stablehlo", "mlir::func::FuncOp"> { |
| 109 | +- let summary = "Legalize TOSA rescales to StableHlo primitive math operations"; |
| 110 | +- let description = [{ |
| 111 | +- This pass rewrites TOSA rescale operations to StableHLO primitive math operations. |
| 112 | +- }]; |
| 113 | +- let dependentDialects = [ |
| 114 | +- "::mlir::stablehlo::StablehloDialect" |
| 115 | +- ]; |
| 116 | +-} |
| 117 | ++// TODO: un-comment the following once #2751 is fixed. |
| 118 | ++// def StablehloQuantLegalizeToTosaRescalePass : Pass<"stablehlo-quant-legalize-to-tosa-rescale", "mlir::func::FuncOp"> { |
| 119 | ++// let summary = "Legalize StableHLO Quantized operations to TOSA rescale operations"; |
| 120 | ++// let description = [{ |
| 121 | ++// This pass rewrites StableHLO quantized operations to integer operations |
| 122 | ++// by inserting TOSA rescale operations at the inputs and outputs of the |
| 123 | ++// integer operations. |
| 124 | ++// }]; |
| 125 | ++// let dependentDialects = [ |
| 126 | ++// "::mlir::tosa::TosaDialect", |
| 127 | ++// ]; |
| 128 | ++// } |
| 129 | ++// |
| 130 | ++// def TosaRescaleLegalizeToStablehloPass : Pass<"tosa-rescale-legalize-to-stablehlo", "mlir::func::FuncOp"> { |
| 131 | ++// let summary = "Legalize TOSA rescales to StableHlo primitive math operations"; |
| 132 | ++// let description = [{ |
| 133 | ++// This pass rewrites TOSA rescale operations to StableHLO primitive math operations. |
| 134 | ++// }]; |
| 135 | ++// let dependentDialects = [ |
| 136 | ++// "::mlir::stablehlo::StablehloDialect" |
| 137 | ++// ]; |
| 138 | ++// } |
1 | 139 | diff --ruN a/stablehlo/stablehlo/dialect/StablehloOps.cpp b/stablehlo/stablehlo/dialect/StablehloOps.cpp |
2 | 140 | --- stablehlo/stablehlo/dialect/StablehloOps.cpp |
3 | 141 | +++ stablehlo/stablehlo/dialect/StablehloOps.cpp |
|
0 commit comments