xai-org/x-algorithm · critical · std::runtime_error

async_emb: ${label} timed out waiting for step ${step}

Error message

async_emb: ${label} timed out waiting for step ${step}

What it means

async_emb's waitLatest blocks until the host pipeline reports completion of an armed step. If hostWaitDone returns false within the deadline, a std::runtime_error is thrown with the pipeline label and step number, meaning the embedding pipeline stalled or the CUDA/NCCL copy never signaled completion.

Source

Thrown at phoenix/xrex/cuda/async_emb/src/async_emb_api.cc:728

  static_assert(
      std::is_invocable_r_v<XLA_FFI_Error*, T, XLA_FFI_CallFrame*>,
      "Encapsulated function must be an XLA FFI handler"
  );
  return nb::capsule(reinterpret_cast<void*>(fn));
}

uint64_t waitLatest(int64_t context_id, AsyncEmbContext::Operation pipeline, const char* label) {
  auto ctx = findContext(context_id);
  if (ctx == nullptr) {
    return 0;
  }
  ctx->ensureHealthy();
  uint64_t step = ctx->armedStep(pipeline);
  if (step == 0) {
    return 0;
  }
  if (!ctx->hostWaitDone(pipeline, step)) {
    throw std::runtime_error(
        "async_emb: " + std::string(label) + " timed out waiting for step " + std::to_string(step)
    );
  }
  return step;
}

nb::bytes testSnapshot(int64_t context_id, const std::string& region) {
  auto ctx = readyContext(context_id);
  if (ctx == nullptr) {
    throw std::invalid_argument("async_emb context not initialized");
  }
  const auto& spec = ctx->spec();
  const auto& layout = ctx->layout();
  const size_t index_bytes =
      size_t(ctx->worldSize()) * size_t(spec.tokens_per_rank) * sizeof(int32_t);
  const size_t block_bytes = size_t(ctx->worldSize()) * size_t(spec.tokens_per_rank) *
                             size_t(spec.shard_width) * sizeof(__nv_bfloat16);
  using Operation = AsyncEmbContext::Operation;

View on GitHub (pinned to 24c60942c5)

Solutions

1. Check GPU health and dmesg/nvidia-smi for Xid errors or hangs; restart the job if a rank died
2. Verify all ranks in the world are alive and participating (a single dead peer stalls collectives)
3. Retry testSnapshot when the pipeline is quiescent rather than mid-step
4. If reproducible, capture NCCL_DEBUG=INFO logs to identify which collective is stuck

# Ensure the pipeline is healthy and a step is armed before snapshotting
step = async_emb.armed_step(ctx_id) if hasattr(async_emb, 'armed_step') else None
if step == 0:
    pass  # nothing in flight; waitLatest returns 0 without blocking

try:
    snap = async_emb._test_snapshot(ctx_id, region)
except RuntimeError as e:
    if "timed out waiting for step" in str(e):
        log.error("pipeline stalled at step, aborting job")
        raise SystemExit(3)  # do not retry a hung collective blindly
    raise

When it happens

Trigger: Calling testSnapshot (which internally calls waitLatest) after the pipeline has armed a step (armedStep != 0) but the host wait times out — e.g. stalled NCCL communicator, dead peer rank, GPU hang, or a snapshot requested while training threads are blocked.

Common situations: One rank crashed or was preempted so collective copies never complete; NCCL timeout/ watchdog misconfiguration; calling testSnapshot during checkpointing or while the pipeline is deliberately paused; driver-level GPU reset.

Understand the failure class

• Timeouts: ETIMEDOUT, deadlines, and hung requests — what actually expires when a request times out.

Related errors

• NCCL ${operation} failed: ${ncclGetErrorString(result)}
• NCCL version query failed: ${ncclGetErrorString(result)}
• could not resolve loaded NCCL library path
• async_emb arena size overflow
• async_emb arena alignment overflow

AI-assisted analysis of xai-org/x-algorithm@24c60942c5 (2026-08-28). Data as JSON: /api/errors/2700d8ff5a72e2c3. Report an issue: GitHub.