}

// CHECK-LABEL: func @add
// CHECK: %[[ADD:.+]] = tfl.add(%arg0, %arg1) <{fused_activation_function = "NONE"}> : (tensor<1x3x!quant.uniform<i8:f32, {{.*}}>>, tensor<1x3x!quant.uniform<i8:f32, {{.*}}>>) -> tensor<1x3x!quant.uniform<i8:f32, {{.*}}>>
// CHECK: return %[[ADD]]

// -----

// Tests that `stablehlo.add` with int32 UniformQuantizedPerAxisTypes is
// not converted.

// RUN: tf-opt %s -tfl-legalize-tf --cse -split-input-file| FileCheck %s --dump-input=fail

func.func @add(%arg0: tensor<1xf32>, %arg1: tensor<1xf32>) -> tensor<1xf32> {
  %0 = "tf.Add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32>
  func.return %0: tensor<1xf32>

// CHECK-LABEL: add
// CHECK:  tfl.add %arg0, %arg1 {fused_activation_function = "NONE"} : tensor<1xf32>
// CHECK:  return
}

		expectedReadyEndpoints: map[types.NamespacedName]int{
			makeNSN("ns1", "ep1"): 2,
			makeNSN("ns2", "ep2"): 1,
		},
	}, {
		// Case[5]: add an Endpoints
		name:              "add an Endpoints",
		previousEndpoints: []*discovery.EndpointSlice{nil},
		currentEndpoints:  namedPortLocal,
		oldEndpoints:      map[proxy.ServicePortName][]endpointExpectation{},

		pendingMap: map[T]struct{}{},
	}
}

func (t *trackingWorkqueue[T]) Add(item T) {
	t.queue(item)
	t.limiter.Add(item)
}
func (t *trackingWorkqueue[T]) AddAfter(item T, duration time.Duration) {
	t.Add(item)
}
func (t *trackingWorkqueue[T]) AddRateLimited(item T) {
	t.Add(item)
}
func (t *trackingWorkqueue[T]) Get() (T, bool) {
	item, shutdown := t.limiter.Get()

int64_t GetConvolutionKernelInputFeatureDimension(bool is_depthwise) {
  return is_depthwise ? 0 : 3;
}

// stablehlo.uniform_quantize -> tfl.quantize
// TODO: b/322428814 - Add StableHLO quantizer integration tests for ODML.
class RewriteUniformQuantizeOp
    : public OpRewritePattern<stablehlo::UniformQuantizeOp> {
  using OpRewritePattern<stablehlo::UniformQuantizeOp>::OpRewritePattern;

		}
		o1 = LOP_IRR(opu, uint32(p.To.Reg), r, uint32(p.From.Offset>>16))
		o2 = LOP_IRR(opl, uint32(p.To.Reg), uint32(p.To.Reg), uint32(p.From.Offset)&0xFFFF)

	case 22: /* add $lcon/$andcon,r1,r2 ==> oris+ori+add/ori+add, add $s34con,r1 ==> addis+ori+slw+ori+add */
		if p.To.Reg == REGTMP || p.Reg == REGTMP {
			c.ctxt.Diag("can't synthesize large constant\n%v", p)
		}
		d := c.vregoff(&p.From)
		r := int(p.Reg)
		if r == 0 {

      "/job:localhost/replica:0/task:0/cpu:0");
  auto x = ops::Placeholder(root.WithOpName("x"), DT_FLOAT);
  auto add1 = ops::Add(root.WithOpName("add1"), x, x);
  add1.node()->AddAttr("_cluster", "cluster1");
  auto add2 = ops::Add(root.WithOpName("add2"), add1, add1);
  add2.node()->AddAttr("_cluster", "cluster2");
  auto out = ops::Mul(root.WithOpName("mul"), add1, add2);

          rewriter.getUnitAttr());
      op.getBiasMutable().assign(none_value);
      return success();
    }

    return failure();
  }
};

// Return true if the given Add operation has the CPU kernel supported shapes.
bool VerifyAddOpShapeConstraints(AddOp op) {
  auto element_type = getElementTypeOrSelf(op.getOutput().getType());

	serialNumber, _ := rand.Int(rand.Reader, serialNumberLimit)

	template := &Certificate{
		SerialNumber: serialNumber,
		Subject:      pkix.Name{CommonName: cn},
		NotBefore:    time.Now().Add(-1 * time.Hour),
		NotAfter:     time.Now().Add(24 * time.Hour),

		KeyUsage:              KeyUsageKeyEncipherment | KeyUsageDigitalSignature | KeyUsageCertSign,
		ExtKeyUsage:           []ExtKeyUsage{ExtKeyUsageServerAuth},

  // expected-remark@above {{ID: 0}}
  // expected-remark@above {{Sinks: {}}}
}

// -----

func.func @add(%arg0: tensor<1xf32>, %arg1: tensor<1xf32>) -> tensor<1xf32> {
  // expected-remark@above {{ID: 2}}
  %sum = "tf.Add"(%arg0, %arg1) : (tensor<1xf32>, tensor<1xf32>) -> tensor<1xf32>
  // expected-remark@above {{ID: 0}}
  func.return %sum : tensor<1xf32>