}
};

// Rewrites quantized `stablehlo.constant` to `tfl.pseudo_qconst`.
class RewriteQuantizedConstantOp
    : public OpRewritePattern<stablehlo::ConstantOp> {
 public:
  using OpRewritePattern<stablehlo::ConstantOp>::OpRewritePattern;

  LogicalResult match(stablehlo::ConstantOp op) const override {
    return success(IsQuantizedTensorType(op.getOutput().getType()));
  }

}

// Return true if the value is a splat tensor constant zero.
bool EqualsZero(Value value) {
  DenseElementsAttr constant;
  if (!matchPattern(value, m_Constant(&constant)) || !constant.isSplat()) {
    return false;
  }

  Type element_type = value.getType().cast<ShapedType>().getElementType();
  if (element_type.isa<FloatType>()) {
    return constant.getSplatValue<APFloat>().isZero();
  } else {

// CHECK-DAG:       [[VAL_92:%.*]] = arith.constant dense<0.000000e+00> : tensor<3xf32>
// CHECK-DAG:       [[VAL_93:%.*]] = arith.constant dense<0.000000e+00> : tensor<1x3xf32>
// CHECK-DAG:       [[VAL_94:%.*]] = arith.constant dense<0.000000e+00> : tensor<1x1xf32>
// CHECK-DAG:       [[VAL_95:%.*]] = arith.constant dense<0> : tensor<1xi64>
// CHECK-DAG:       [[VAL_96:%.*]] = arith.constant dense<1> : tensor<1xi64>

// CHECK-DAG:       [[VAL_8:%.*]] = arith.constant dense<0.000000e+00> : tensor<16x16xf32>
// CHECK-DAG:       [[VAL_9:%.*]] = arith.constant dense<0.000000e+00> : tensor<16x8xf32>
// CHECK-DAG:       [[VAL_10:%.*]] = arith.constant dense<0.000000e+00> : tensor<16xf32>
// CHECK-DAG:       [[VAL_11:%.*]] = arith.constant dense<0.000000e+00> : tensor<1x16xf32>
// CHECK-DAG:       [[VAL_12:%.*]] = arith.constant dense<0.000000e+00> : tensor<8x16xf32>

// CHECK-DAG: %[[BROADCAST_DIM:.+]] = arith.constant dense<{{\[3, 2, 1, 1, 4\]}}> : tensor<5xi32>
// CHECK-DAG: %[[EXPAND_DIM1:.+]] = arith.constant dense<3> : tensor<1xi32>
// CHECK-DAG: %[[EXPAND_DIM0:.+]] = arith.constant dense<2> : tensor<1xi32>
// CHECK-DAG: %[[PERM:.+]] = arith.constant dense<{{\[1, 0, 2\]}}> : tensor<3xi32>

literal.
</p>


<h2 id="Constants">Constants</h2>

<p>There are <i>boolean constants</i>,
<i>rune constants</i>,
<i>integer constants</i>,
<i>floating-point constants</i>, <i>complex constants</i>,
and <i>string constants</i>. Rune, integer, floating-point,
and complex constants are
collectively called <i>numeric constants</i>.
</p>

<p>
A constant value is represented by a

}

static bool IsConst(Operation* op) {
  return isa<mlir::func::ConstantOp, mlir::arith::ConstantOp, mlir::TF::ConstOp,
             tfl::ConstOp, tfl::QConstOp, tfl::SparseConstOp,
             tfl::SparseQConstOp, mlir::TFL::NoValueOp,
             mlir::stablehlo::ConstantOp, mlir::vhlo::ConstantOpV1>(op);
}

static bool IsTFResourceOp(Operation* op) {

	NORMAL_PRIORITY_CLASS         = 0x00000020
	PROCESS_MODE_BACKGROUND_BEGIN = 0x00100000
	PROCESS_MODE_BACKGROUND_END   = 0x00200000
	REALTIME_PRIORITY_CLASS       = 0x00000100
)

/* wintrust.h constants for WinVerifyTrustEx */
const (
	WTD_UI_ALL    = 1
	WTD_UI_NONE   = 2
	WTD_UI_NOBAD  = 3
	WTD_UI_NOGOOD = 4

	WTD_REVOKE_NONE       = 0
	WTD_REVOKE_WHOLECHAIN = 1

	WTD_CHOICE_FILE    = 1

		return C_SPR
	}
	return C_GOK
}

// con32class reclassifies the constant of 32-bit instruction. Because the constant type is 32-bit,
// but saved in Offset which type is int64, con32class treats it as uint32 type and reclassifies it.
func (c *ctxt7) con32class(a *obj.Addr) int {
	v := uint32(a.Offset)
	// For 32-bit instruction with constant, rewrite
	// the high 32-bit to be a repetition of the low

				}},
			},
			"", "",
		},
	}
	for _, tt := range tests {
		t.Run(tt.name, func(t *testing.T) {
			annotations := map[string]string{}
			annotations[constants.InternalGatewaySemantics] = constants.GatewaySemanticsGateway

			warn, err := ValidateGateway(config.Config{
				Meta: config.Meta{
					Name:        someName,
					Namespace:   someNamespace,
					Annotations: annotations,