op)))
        return failure();

      OpBuilder builder(op);
      auto const_op = builder.create<TF::ConstOp>(
          op->getLoc(), DenseIntElementsAttr::get(
                            RankedTensorType::get({}, builder.getI64Type()),
                            {device_ordinal}));
      op->replaceAllUsesWith(const_op);
      op->erase();
      return WalkResult::advance();
    }

      ModuleOp::create(mlir::UnknownLoc::get(entry_func.getContext()),
                       absl::StrCat("module_", module_name));
  auto parent_module = entry_func->getParentOfType<ModuleOp>();
  auto versions_attr = parent_module->getAttr(kVersionsAttr);
  if (!versions_attr)
    return parent_module.emitError(CreateMissingAttributeMsg(kVersionsAttr));

  module_for_func.get().getOperation()->setAttr(kVersionsAttr, versions_attr);

    // types are static.
    SmallVector<DenseIntElementsAttr, 4> indices;
    indices.reserve(op.getN());
    for (auto it : llvm::zip(op.getIndices(), op.getData())) {
      Value index = std::get<0>(it);
      Value data = std::get<1>(it);

      DenseIntElementsAttr index_attr;
      if (!matchPattern(index, m_Constant(&index_attr))) return failure();

  if (!HasRankAtLeast(op.getX(), 2)) {
    return op.emitOpError("requires lhs operand to have rank at least two");
  }
  if (!HasRankAtLeast(op.getY(), 2)) {
    return op.emitOpError("requires rhs operand to have rank at least two");
  }

  RankedTensorType x_ty = GetRankedTensorTypeForOperand(op.getX());
  RankedTensorType y_ty = GetRankedTensorTypeForOperand(op.getY());

  if (!x_ty || !y_ty) return success();

    // `stablehlo.recv`.
    effects.emplace_back(MemoryEffects::Write::get(),
                         ResourceEffects::Send::get());
    effects.emplace_back(MemoryEffects::Write::get(),
                         ResourceEffects::Recv::get());
    effects.emplace_back(MemoryEffects::Write::get(),
                         ResourceEffects::XlaHostCompute::get());
  }
}

LogicalResult XlaCallModuleOp::verifySymbolUses(

        op->getLoc(),
        RankedTensorType::get(permuted_shape, input_type.getElementType()),
        input, permutation_tensor_op.getResult());
  };

  if (!can_convert_to_fully_connected(op)) {
    return false;
  }

  Value input_lhs = get_real_input_value(op.getX());
  Value input_rhs = get_real_input_value(op.getY());

  Value legalized_lhs =

            interface A {
                String getName()
                void setName(String name)

                B getB()
            }

            @Managed
            interface B {
                C getC()
            }

            @Managed
            interface C {
                A getA()
                void setA(A a)
            }

            class RulePlugin extends RuleSource {

    @Managed
    interface A1 {
        A1 getA();

        B1 getB();

        C1 getC();

        D1 getD();
    }

    @Managed
    interface B1 {
        A1 getA();

        B1 getB();

        C1 getC();

        D1 getD();
    }

    @Managed
    interface C1 {
        A1 getA();

        B1 getB();

        C1 getC();

		//	38	eattr[4]
		//	42	off[4]
		//	46	name[namelen]
		//	46+namelen+xlen+fclen - next header
		//
		if get4(buf) != zcheader {
			break
		}
		meth := get2(buf[10:])
		size := get4(buf[24:])
		namelen := get2(buf[28:])
		xlen := get2(buf[30:])
		fclen := get2(buf[32:])
		off := get4(buf[42:])
		zname := buf[46 : 46+namelen]
		buf = buf[46+namelen+xlen+fclen:]
		if string(zname) != name {
			continue

func (b *wbBuf) empty() bool {
	return b.next == uintptr(unsafe.Pointer(&b.buf[0]))
}

// getX returns space in the write barrier buffer to store X pointers.
// getX will flush the buffer if necessary. Callers should use this as:
//
//	buf := &getg().m.p.ptr().wbBuf
//	p := buf.get2()
//	p[0], p[1] = old, new
//	... actual memory write ...
//