auto* devce_hardware = GetTargetHardware(hardware);
  if (devce_hardware == nullptr) return {};
  return devce_hardware->IsOpSupported(op);
}

// ================== Convert Quantized Op ============================

// Walk through the func and convert the quantize ops to their float version.
void ConvertQuantizedOpToFloat(mlir::func::FuncOp func, OpBuilder* builder) {
  func.walk([&](Operation* op) {

				// could have been deleted, or could be no items in queue not currently worked on.  need a way to differentiate.
				wp.lock.Unlock()
				continue
			}
			wp.q.Delete(target)
			wp.currentlyWorking.Insert(convert(target))
			wp.lock.Unlock()
			// work should be done without holding the lock
			cfg := wp.get(target)
			if cfg != nil {
				// Check that generation matches

#include "stablehlo/dialect/StablehloOps.h"  // from @stablehlo
#include "tensorflow/compiler/mlir/tensorflow/ir/tf_ops.h"

namespace mlir {
namespace odml {

namespace {

// Convert op to stablehlo.custom_call
//   "tf.ResizeBilinear"(%341, %138) {
//      align_corners = false, device = "", half_pixel_centers = true}
//   ==>
//   stablehlo.custom_call @tf.ResizeBilinear(%arg0, %arg1) {

			// create a match all object
			cpw.Match = &networking.EnvoyFilter_EnvoyConfigObjectMatch{Context: networking.EnvoyFilter_ANY}
		} else if cp.Match.Proxy != nil && cp.Match.Proxy.ProxyVersion != "" {
			// Attempt to convert regex to a simple prefix match for the common case of matching
			// a standard Istio version. This field should likely be replaced with semver, but for now
			// we can workaround the performance impact of regex

		[6]*argField{ap_VecReg_6_10, ap_VecReg_16_20}},
	{XSCVSDQP, 0xfc1f07fe00000000, 0xfc0a068800000000, 0x100000000, // VSX Scalar Convert Signed Doubleword to Quad-Precision format X-form (xscvsdqp VRT,VRB)
		[6]*argField{ap_VecReg_6_10, ap_VecReg_16_20}},

// RUN: tf-quant-opt %s -quant-convert-fake-quant-to-qdq -quant-lift-quantizable-spots-as-functions -quant-insert-quantized-functions -quant-quantize-composite-functions -symbol-dce | FileCheck %s

func.func @fake_quant_conv(%arg0: tensor<1x3x4x3xf32>, %arg1: tensor<2x3x3x2xf32>) -> tensor<*xf32> {
  %cst = "tf.Const"() {value = dense<0.000000e+00> : tensor<2xf32>} : () -> tensor<2xf32>

						return true, proto.Clone(lp.Value).(*listener.ListenerFilter)
					}
					return false, nil
				},
			)
		} else if lp.Operation == networking.EnvoyFilter_Patch_INSERT_BEFORE {
			// insert before without a filter match is same as insert in the beginning
			if !hasListenerFilterMatch(lp) {
				lis.ListenerFilters = append([]*listener.ListenerFilter{proto.Clone(lp.Value).(*listener.ListenerFilter)}, lis.ListenerFilters...)
				continue

      const bool patch =
          (m.op_index >= remat.insert) && (m.op_index < remat.insert + len);
      // Are we past the insertion portion of the new operation sequence?
      // Then we need to convert indices back to the original sequence.
      const int shift = (m.op_index >= remat.insert + len) ? len : 0;
      m.size += patch ? 0 : operations_[m.op_index - shift].alloc;

// folding logic will use a "arith.constant" op to replace the
// "tf.FakeQuantWithMinMaxVarsOp", the "tfl.quantize" op is used to preserve
// the quantization parameters as a TypeAttr and "tfl.dequantize" op used to
// convert the output type to the next op. Here are the transformations:
//
// input   min cst       max cst          input   min cst       max cst
//  \       |             |                \       |             |

  LegalizeTFToQuant(const LegalizeTFToQuant &) {}

  /// Performs the lowering to Quant ops dialect.
  void runOnOperation() override;

  void getDependentDialects(DialectRegistry &registry) const override {
    registry.insert<quant::QuantizationDialect,
                    quantfork::QuantizationForkDialect>();
  }

  StringRef getArgument() const final {
    // This is the argument used to refer to the pass in