//go:nosplit
func (i *Int32) Swap(new int32) int32 {
	return Xchgint32(&i.value, new)
}

// Add adds delta to i atomically, returning
// the new updated value.
//
// This operation wraps around in the usual
// two's-complement way.
//
//go:nosplit
func (i *Int32) Add(delta int32) int32 {
	return Xaddint32(&i.value, delta)
}

// Int64 is an atomically accessed int64 value.
//

  CellReader<int64_t> static_error(kNonStaticOpStreamz);
  CreateModule(kNonStaticFailure);

  auto result = Run();

  EXPECT_TRUE(result.failed());
  EXPECT_EQ(legal_error.Delta("mhlo.dynamic_iota"), 0);
  EXPECT_EQ(static_error.Delta("mhlo.dynamic_iota"), 1);
}

TEST_F(VerifyTfxlaLegalizationTest, SkipsSpecificNonStaticInputs) {
  // Using a string constant here instead of testdata to make this compatible

}

func (v *Float) appendJSON(b []byte) []byte {
	return strconv.AppendFloat(b, math.Float64frombits(v.f.Load()), 'g', -1, 64)
}

// Add adds delta to v.
func (v *Float) Add(delta float64) {
	for {
		cur := v.f.Load()
		curVal := math.Float64frombits(cur)
		nxtVal := curVal + delta
		nxt := math.Float64bits(nxtVal)
		if v.f.CompareAndSwap(cur, nxt) {
			return
		}
	}
}

// Set sets v to value.

          <y:Arrows source="delta" target="none"/>
        </y:QuadCurveEdge>
      </data>
    </edge>
    <edge id="e1" source="n5" target="n4">
      <data key="d9">
        <y:QuadCurveEdge straightness="0.1">
          <y:Path sx="70.0" sy="-0.0" tx="-70.0" ty="-10.0"/>
          <y:LineStyle color="#000000" type="line" width="1.0"/>
          <y:Arrows source="delta" target="none"/>
        </y:QuadCurveEdge>

	return image.DecodeConfig(bufio.NewReader(f))
}

func delta(u0, u1 uint32) int {
	d := int(u0) - int(u1)
	if d < 0 {
		return -d
	}
	return d
}

func withinTolerance(c0, c1 color.Color, tolerance int) bool {
	r0, g0, b0, a0 := c0.RGBA()
	r1, g1, b1, a1 := c1.RGBA()
	r := delta(r0, r1)
	g := delta(g0, g1)
	b := delta(b0, b1)
	a := delta(a0, a1)

  })";

  CellReader<int64_t> error(kHasTpuPartitionedCallStreamzName);
  CreateModule(kMlirModuleStr);

  auto result = Run();

  EXPECT_TRUE(result.succeeded());

  EXPECT_EQ(error.Delta("true"), 1);
  EXPECT_EQ(error.Delta("false"), 0);
}

TEST_F(InferenceMetricsPassTest, RecordsFalseForNonTPUPartitionedCallOp) {
  static constexpr char kMlirModuleStr[] = R"(

	ScrapeTypeEnvoy = "envoy"
	ScrapeTypeApp   = "application"
	ScrapeTypeAgent = "agent"
)

var processStartTime = time.Now()

func RecordStartupTime() {
	delta := time.Since(processStartTime)
	startupTime.Record(delta.Seconds())
	log.Infof("Readiness succeeded in %v", delta)

  }

  /**
   * Adds {@code delta} to the value currently associated with {@code key}, and returns the new
   * value.
   */
  @CanIgnoreReturnValue
  public long addAndGet(K key, long delta) {
    outer:
    while (true) {
      AtomicLong atomic = map.get(key);
      if (atomic == null) {
        atomic = map.putIfAbsent(key, new AtomicLong(delta));
        if (atomic == null) {

  EXPECT_TRUE(result.ok());
  EXPECT_EQ(compilation_status.Delta(kMlirCombinedMlirSuccess), 1);
  EXPECT_EQ(compilation_status.Delta(kMlirCombinedMlirFailure), 0);
  EXPECT_EQ(compilation_status.Delta(kMlirCombinedOldSuccess), 1);
  EXPECT_EQ(compilation_status.Delta(kMlirCombinedOldFailure), 0);
  // Old bridge should never be called at all.
  EXPECT_EQ(compilation_status.Delta(kOldBridgeMlirFilteredFailure), 0);

          <y:Arrows source="delta" target="none"/>
        </y:QuadCurveEdge>
      </data>
    </edge>
    <edge id="e1" source="n3" target="n6">
      <data key="d9">
        <y:QuadCurveEdge straightness="0.1">
          <y:Path sx="70.0" sy="7.5" tx="-70.0" ty="-0.0"/>
          <y:LineStyle color="#000000" type="line" width="1.0"/>
          <y:Arrows source="delta" target="none"/>
        </y:QuadCurveEdge>