// CHECK-SAME: channel_handle = #mhlo.channel_handle
  // CHECK-SAME: handle =
  // CHECK-SAME: type = 2
  // CHECK-SAME: mhlo.frontend_attributes = {_xla_host_transfer_handler_name = "tf_rendezvous", _xla_host_transfer_rendezvous = "send0_dtoh_0"}

  // CHECK:      [[RECV0_RETVAL0_TUPLE:%.*]]:2 = "mhlo.recv"([[SEND0_ARG0_TOKEN]])
  // CHECK-SAME: channel_handle = #mhlo.channel_handle
  // CHECK-SAME: handle =

  %group_assignment = "tf.Const"() { value = dense<[[0],[1]]> : tensor<2x1xi32> } : () -> tensor<2x1xi32>
  // CHECK: "mhlo.all_reduce"
  // CHECK-SAME: channel_handle = #mhlo.channel_handle<handle = 10001, type = 1>
  // CHECK-SAME{LITERAL}: replica_groups = dense<[[0], [1]]> : tensor<2x1xi64>
  // CHECK: mhlo.add
  // CHECK: mhlo.return

  builder.setInsertionPoint(op);
  ChannelHandleAttr channel_handle = ConvertChannel(builder, channel_id, mode);
  Location loc = op->getLoc();
  Type element_type = getElementTypeOrSelf(input.getType());
  auto all_reduce = builder.create<AllReduceOp>(
      loc, result_type, input, replica_groups, channel_handle, nullptr);

  if (all_reduce.getNumResults() != 1) {
    return op->emitOpError()

  // type 2 == DEVICE_TO_HOST
  auto channel_handle = ChannelHandleAttr::get(builder.getContext(),
                                               /*handle=*/GetNextChannelId(),
                                               /*type=*/2);
  auto send = builder.create<SendOp>(
      loc, token.getType(), operand, token, channel_handle,
      /*is_host_transfer=*/builder.getBoolAttr(true));