// typexpr    <expr> (               <mode>                    )
//
// constant   <expr> (<untyped kind> <mode>                    )
// constant   <expr> (               <mode>       of type <typ>)
// constant   <expr> (<untyped kind> <mode> <val>              )
// constant   <expr> (               <mode> <val> of type <typ>)
//
// variable   <expr> (<untyped kind> <mode>                    )

// typexpr    <expr> (               <mode>                    )
//
// constant   <expr> (<untyped kind> <mode>                    )
// constant   <expr> (               <mode>       of type <typ>)
// constant   <expr> (<untyped kind> <mode> <val>              )
// constant   <expr> (               <mode> <val> of type <typ>)
//
// variable   <expr> (<untyped kind> <mode>                    )

  func.return %rst : tensor<256x8x7x16xf32>

// CHECK-DAG: %[[CONSTANT:.*]] = arith.constant dense<0.000000e+00> : tensor<16xf32>
// CHECK-DAG: %[[CONSTANT0:.*]] = arith.constant dense<0.000000e+00> : tensor<16x3x3x3xf32>
// CHECK: %[[QUANTIZE:.*]] = "tfl.quantize"(%[[CONSTANT0]]) <{qtype = tensor<16x3x3x3x!quant.uniform<u4:f32, 1.000000e+00>>}>

  func.return %rst : tensor<256x8x7x16xf32>

// CHECK-DAG: %[[CONSTANT:.*]] = arith.constant dense<0.000000e+00> : tensor<16xf32>
// CHECK-DAG: %[[CONSTANT0:.*]] = arith.constant dense<0.000000e+00> : tensor<16x3x3x3xf32>
// CHECK: %[[QUANTIZE:.*]] = "tfl.quantize"(%[[CONSTANT0]]) <{qtype = tensor<16x3x3x3x!quant.uniform<u8:f32, 1.000000e+00>>}>

  return %0 : tensor<16x4xf32>
}
// Check that the constant is cloned with same value.
// The duplicated constant is the last index of the ConcatV2 op (which
// accepts a variadic arg).
// CHECK-DAG: %[[CST_0:.*]] = "tf.Const"()
// CHECK-SAME: value = dense<1> : tensor<i32>
// CHECK-DAG: %[[CST_1:.*]] = "tf.Const"()
// CHECK-SAME: value = dense<1> : tensor<i32>

// Check that the constants used for tf.ConcatV2 and tf.AddV2 are different.

// function signature. Constants are also transposed accordingly.

// CHECK-LABEL: nchw_conv_with_bias_add
// CHECK-SAME: %[[ARG:.+]]: tensor<1x2x5x5xf32>
func.func @nchw_conv_with_bias_add(%arg0: tensor<1x2x5x5xf32>) -> tensor<1x4x5x5xf32> {
  %0 = stablehlo.constant dense<2.000000e+00> : tensor<4x2x3x3xf32>
  %1 = stablehlo.constant dense<3.000000e+00> : tensor<1x4x5x5xf32>

// only be used as the shape-determining operand; it will not replace other
// usages of the original constant. If the operands are not constants (i.e.
// results of some other computation), then the pass recursively traverses the
// call tree upwards and duplicates all constants found in the subtree in a
// similar manner.
//
// This pass may be used to avoid placing shape-determining constants in the CPU

// for debugging purposes.
func (po *poset) CheckIntegrity() {
	// Record which index is a constant
	constants := newBitset(int(po.lastidx + 1))
	for _, c := range po.constants {
		constants.Set(c)
	}

	// Verify that each node appears in a single DAG, and that
	// all constants are within the first DAG
	seen := newBitset(int(po.lastidx + 1))
	for ridx, r := range po.roots {
		if r == 0 {

	VPMSUMD	V19,const1,V11	// vpmsumd with constants
	LVX	(R4+off48),V19	// load next from buffer
	OR	$0,R2,R2

	VPMSUMD	V20,const1,V12	// vpmsumd with constants
	LVX	(R4+off64),V20	// load next from buffer
	OR	$0,R2,R2

	VPMSUMD	V21,const1,V13	// vpmsumd with constants
	LVX	(R4+off80),V21	// load next from buffer
	OR	$0,R2,R2

	VPMSUMD	V22,const1,V14	// vpmsumd with constants
	LVX	(R4+off96),V22	// load next from buffer

	switch typ := under(x.typ).(type) {
	case *Basic:
		if isString(typ) {
			valid = true
			if x.mode == constant_ {
				length = int64(len(constant.StringVal(x.val)))
			}
			// an indexed string always yields a byte value
			// (not a constant) even if the string and the
			// index are constant
			x.mode = value
			x.typ = universeByte // use 'byte' name
		}

	case *Array:
		valid = true
		length = typ.len