serve: a checkpoint save that fails with a device fault reports no reason to the client, and gpu_sync_ok() passed on the already-faulted device (240K prefill, 0.1.41, gfx1100) #1932

Environment

Relevant args: --max-context 262144 --kv int8 --kv-resident 32768 --spec 4 --prefill auto:32768 --expert-cache 4096 --vram-reserve-mib 700, plus a CPU vision encoder.

What happened

A single request with a 240,903-token prompt died part way through prefill with a 400 the client cannot explain, and the engine had to reload the model:

[strata] reading the prompt: 194,560 of 240,903 tokens, 192 s so far
[strata] the engine reported an error: saving a conversation checkpoint failed
[strata] done: 0 tokens in 192 s (0.0 tok/s) (error, cancel=False)
[strata] 400 invalid request: saving a conversation checkpoint failed
[strata] model unloaded; the next request loads it again

The cause is in the engine log, one line above the summary:

strata serve: prompt chunk auto: 19456 tokens, a 384-slot ring
strata serve: the prompt path borrows 5177 CUDA0 cache slots (8.39 GiB)
strata serve: 2673 MiB of VRAM free with everything loaded
strata serve: checkpoint save: conversation snapshot running-state copy: unspecified launch failure
strata serve: saving a conversation checkpoint failed
hipModuleUnload failed:
 error: unspecified launch failure
hipModuleUnload failed:
 error: unspecified launch failure
hipModuleUnload failed:
 error: unspecified launch failure

Note where it failed: 194560 = 10 x 19456 — exactly the end of the 10th prompt chunk, i.e. the periodic checkpoint save. 2673 MiB of VRAM was free, so this is not VRAM exhaustion.

I have seen this once (0.1.41, 2026-10-11). I am not able to reproduce it on demand; many other 240K prefills on the same machine and config went through fine, before and after.

Two things that look like real defects, independent of the fault itself

1. gpu_sync_ok() passes on an already-faulted device, so the message misattributes the failure.

src/program/generate.cpp:7217-7223:

// #613: why the last checkpoint_at failed - a failed device sync is the GPU itself (a hang Windows then resets),
// not the checkpoint, and the error says so
std::string ckpt_why;
auto gpu_sync_ok = [&]() -> bool {
    const cudaError_t e = cudaDeviceSynchronize();
    if (e == cudaSuccess) return true;
    ckpt_why = std::string(": the GPU stopped responding before it (") + cudaGetErrorString(e) +
               ") - a GPU hang; on Windows the driver is then reset";
    return false;
};

Here cudaDeviceSynchronize() succeeded — ckpt_why stayed empty, which is why the request's error carries no reason — and the very next operation failed:

if (!gpu_sync_ok() || !checkpoint_save(c, ss, g)) return false;   // generate.cpp:7247

checkpoint_save reported conversation snapshot running-state copy: unspecified launch failure, i.e. a device fault (hipErrorLaunchFailure) on the copy. The three hipModuleUnload failed lines immediately after confirm the context really was dead. So the #613 discrimination did not fire on a case that was, in fact, the GPU — it just reported a bare "checkpoint failed", which points the reader at the checkpoint code instead.

2. The reason never reaches the client, only stderr.

src/program/generate.cpp:1380-1385:

bool checkpoint_save(ConvCheckpoint& c, const SessionState& ss, const ModelGeometry& g) {
    std::string error;
    if (strata::core::conversation_checkpoint_save(c, ss, g, error)) return true;
    std::fprintf(stderr, "strata serve: checkpoint save: %s\n", error.c_str());   // the caller's ERR has no reason
    return false;
}

The comment acknowledges the gap. From an API client's point of view a failed checkpoint save is indistinguishable from any other 400, and since the request took 192 s to get there, there is nothing to go on. I only found the reason because this machine keeps the engine's stderr in a log file. Would it be acceptable to append ckpt_why (or the error string) to the caller's e unconditionally, the way the GPU-hang path already does?

Question rather than a claim

Is aborting the whole request the intended response to a failed checkpoint save? A checkpoint is only a cache for the next continuation, so in principle the current prefill could finish and merely lose reuse. Here it may be moot — the context was already faulted and the engine reloaded anyway — but if checkpoint_at() can also fail for benign reasons, discarding a 240K prefill (192 s) over a failed cache save is a harsh trade.

I have no evidence the checkpoint save caused the fault; it may simply be the first device operation after it. If it is any help, the common factor with #613 is a 240K-class prefill on Windows + gfx1100 with the cache lent to the prompt path (the prompt path borrows 5177 CUDA0 cache slots), and the failure landed exactly on a chunk boundary where the save happens.