}
    rewriter.replaceOp(op, ValueRange(results));
    return success();
  }
};

// Converts tf.OutfeedEnqueueTuple to XLA HLO tuple, create_token and outfeed
// ops.
//
// XLA HLO outfeed op expects a token, which we generate by emitting an
// create_token op.
//
// For example the following IR:
// "tf.OutfeedEnqueueTuple"(%val_1, %val_2) : (tensor<3xi32>, tensor<4xf32>) ->

    TF_VariantTensor:$optional
  );
}

def TF_OutfeedEnqueueOp : TF_Op<"OutfeedEnqueue", []> {
  let summary = "Enqueue a Tensor on the computation outfeed.";

  let arguments = (ins
    Arg<TF_Tensor, [{A tensor that will be inserted into the outfeed queue.}]>:$input
  );

  let results = (outs);

  TF_DerivedOperandTypeAttr dtype = TF_DerivedOperandTypeAttr<0>;
}

func.func @outfeed_enqueue_tuple(%data_1: tensor<3xi32>, %data_2: tensor<4xf32>) -> () {
// CHECK: [[TOKEN:%.*]] = mhlo.create_token  : !mhlo.token
// CHECK: "mhlo.outfeed"([[VAL_0]], [[VAL_1]], [[TOKEN]]) <{outfeed_config = ""}> : (tensor<3xi32>, tensor<4xf32>, !mhlo.token) -> !mhlo.token
  "tf.OutfeedEnqueueTuple"(%data_1, %data_2) : (tensor<3xi32>, tensor<4xf32>) -> ()
  func.return
}

	if err == nil {
		fd[0] = int(fdx[0])
		fd[1] = int(fdx[1])
	}
	return
}

func Sendfile(outfd int, infd int, offset *int64, count int) (written int, err error) {
	if race.Enabled {
		race.ReleaseMerge(unsafe.Pointer(&ioSync))
	}
	return sendfile(outfd, infd, offset, count)
}

}

func Mknod(path string, mode uint32, dev uint64) (err error) {
	return Mknodat(AT_FDCWD, path, mode, dev)
}

func Sendfile(outfd int, infd int, offset *int64, count int) (written int, err error) {
	if raceenabled {
		raceReleaseMerge(unsafe.Pointer(&ioSync))
	}
	return sendfile(outfd, infd, offset, count)
}

//sys	ptrace(request int, pid int, addr uintptr, data int) (err error)

#define ZERO    V0
#define IN0     V1
#define IN1     V2
#define KEY     V3
#define RCON    V4
#define MASK    V5
#define TMP     V6
#define STAGE   V7
#define OUTPERM V8
#define OUTMASK V9
#define OUTHEAD V10
#define OUTTAIL V11

// For P9 instruction emulation
#define ESPERM  V21  // Endian swapping permute into BE
#define TMP2    V22  // Temporary for P8_STXVB16X/P8_STXVB16X

// For {en,de}cryptBlockAsm

	}
	return reclen - uint64(unsafe.Offsetof(Dirent{}.Name)), true
}

func Gettimeofday(tv *Timeval) (err error) {
	err = gettimeofday(tv, nil)
	return
}

// TODO
func sendfile(outfd int, infd int, offset *int64, count int) (written int, err error) {
	return -1, ENOSYS
}

//sys	getdirent(fd int, buf []byte) (n int, err error)

func ReadDirent(fd int, buf []byte) (n int, err error) {

	return mountSetattr(dirfd, pathname, flags, attr, unsafe.Sizeof(*attr))
}

func Sendfile(outfd int, infd int, offset *int64, count int) (written int, err error) {
	if raceenabled {
		raceReleaseMerge(unsafe.Pointer(&ioSync))
	}
	return sendfile(outfd, infd, offset, count)
}

// Sendto
// Recvfrom
// Socketpair

/*
 * Direct access
 */

            preset_method=_PresetMethod.METHOD_STATIC_RANGE_INT8
        ),
        tags=tags,
        signature_keys=['serving_default'],
        op_set=quant_opts_pb2.XLA,
        # Gather op is opt-outed if the size is smaller than the threshold.
        min_num_elements_for_weights=1024 if expect_quantized_gather else 8192,
    )

    data_gen = self._create_data_generator(
        input_key='input_tensor',