Node* recv = RecvAtHost(
        ops::NodeOut(key_constant, 0), "F1", "F1", "O1", {DT_FLOAT, DT_FLOAT},
        b2.opts().WithAttr(kXlaHasHostTransferAttrName, true));
    Node* e = Binary(ops::NodeOut(recv, 0), ops::NodeOut(recv, 1),
                     b2.opts()
                         .WithName("E")
                         .WithControlInputs({recv})

			}
			if !cas.recv.IsValid() {
				t.Fatalf("%s\nselected #%d but internal error: missing recv value", fmtSelect(info), i)
			}
			if recv.Interface() != cas.recv.Interface() || recvOK != !cas.closed {
				if recv.Interface() == cas.recv.Interface() && recvOK == !cas.closed {

	return v.recv(false)
}

// internal recv, possibly non-blocking (nb).
// v is known to be a channel.
func (v Value) recv(nb bool) (val Value, ok bool) {
	tt := (*chanType)(unsafe.Pointer(v.typ()))
	if ChanDir(tt.Dir)&RecvDir == 0 {
		panic("reflect: recv on send-only channel")
	}
	t := tt.Elem
	val = Value{t, nil, flag(t.Kind())}
	var p unsafe.Pointer
	if t.IfaceIndir() {
		p = unsafe_New(t)
		val.ptr = p

    Graph* g, const string& oc_cluster_name,
    const std::vector<DataType>& recv_at_host_dtypes, Node* key_placeholder) {
  NodeDefBuilder recv_at_host_builder(
      absl::StrCat("outside_compilation_", oc_cluster_name, "_recv"),
      "_XlaRecvAtHost");
  NodeDef recv_at_host_def;
  recv_at_host_builder.Attr("Toutputs", recv_at_host_dtypes);
  // The correct device_ordinal will be inserted during replication in a
  // subsequent rewrite.

	"⨔":                         "\u2a14",
	"⊀":                             "\u2280",
	"⋠":                          "\u22e0",
	"⪯̸":                            "\u2aaf\u0338",
	"⊀":                           "\u2280",
	"⪯̸":                         "\u2aaf\u0338",
	"⇏":                           "\u21cf",
	"↛":                           "\u219b",

  let constructor = "TFTPU::CreateTPUParallelExecuteSinkResourceWritePass()";
}

def RewriteTPUEmbeddingOpsPass : Pass<"tf-rewrite-tpu-embedding-ops", "mlir::func::FuncOp"> {
  let summary = "Rewrites TPU embedding send/recv ops by adding TPU embedding "
           "deduplication data";

  let constructor = "TF::CreateRewriteTPUEmbeddingOpsPass()";
}

		"nparallel;":                       '\U00002226',
		"npolint;":                         '\U00002A14',
		"npr;":                             '\U00002280',
		"nprcue;":                          '\U000022E0',
		"nprec;":                           '\U00002280',
		"nrArr;":                           '\U000021CF',
		"nrarr;":                           '\U0000219B',
		"nrightarrow;":                     '\U0000219B',

    // `stablehlo.custom_call` will be lowered to `stablehlo.send` and
    // `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(),

						// in this window, back off to give pp a chance to
						// schedule runnext. This will avoid thrashing gs
						// between different Ps.
						// A sync chan send/recv takes ~50ns as of time of
						// writing, so 3us gives ~50x overshoot.
						if !osHasLowResTimer {
							usleep(3)
						} else {
							// On some platforms system timer granularity is