server: preemption notices, asynchronous parks and exact concurrency (pin branch on b10869) - #210
Conversation
Codex Review SummaryThis comment shows the latest Codex review activity on this pull request.
ℹ️ About Codex in GitHubYour team has set up Codex to review pull requests in this repo. Reviews are triggered when you
Codex reacts with 👀 while any review is running, comments if it has suggestions, and reacts with 👍 once all reviews finish with no findings. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: afd3248b82
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| echo "| delete failures | $failed |" | ||
| } >> "$GITHUB_STEP_SUMMARY" | ||
|
|
||
| # ── Keep the ccache budget inside the repo limit ── |
There was a problem hiding this comment.
This newly added comment uses Unicode box-drawing characters, while the repository standard explicitly prohibits Unicode characters and requires ASCII equivalents. Replace the separators with ordinary hyphens.
AGENTS.md reference: AGENTS.md:L69-L73
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
Moot after 62cf502: this file was a stale copy of the fork's own workflow carried in from the branch point and is no longer in the pin tree. See the reply on the matrix-waiter thread.
| # Everything in the run except this job and the alert job that | ||
| # reports on it. `resolve` stays in the set; it is already a | ||
| # `needs:`, so it is a free consistency check. | ||
| SIBS="$(printf '%s\n' "$JOBS" | awk -F'\t' -v self="$SELF" -v alert="$ALERT" 'NF && $1 != self && $1 != alert')" |
There was a problem hiding this comment.
The jobs API also returns queued jobs, but this filter excludes only the assemble and alert jobs. The reclaim job later declares needs: [..., assemble], so it cannot complete until this job does; once its queued record appears in SIBS, PENDING can never reach zero and every build that reaches assemble waits until the 330-minute deadline instead of publishing. Filter out Reclaim artifact storage as well, or restrict SIBS to the six build prefixes plus resolve.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
These workflow files were stale copies of the fork's own nightly workflows carried in from the branch point, not part of the pin. 62cf502 drops all 40 fork-owned files (unsloth-*.yml, the prebuilt-alert action, scripts/unsloth) from the pin tree; the nightly composes those from master. The three items on those files are therefore moot here; if the matrix waiter or the ARM label are wrong on master they belong in a master PR.
| matrix: | ||
| include: | ||
| - { arch: x64, runner: ubuntu-22.04 } | ||
| - { arch: arm64, runner: ubuntu-22.04-arm } |
There was a problem hiding this comment.
The parent workflow explicitly records that GitHub-hosted arm64 is available only as ubuntu-24.04-arm and uses that label for CUDA, so this ubuntu-22.04-arm leg cannot be scheduled; the Vulkan workflow repeats the invalid label at line 57. Since full publication requires both ARM artifacts, these jobs block every release. Point both jobs at an available ARM runner or provision a self-hosted Jammy runner with this label.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
Moot after 62cf502: this file was a stale copy of the fork's own workflow carried in from the branch point and is no longer in the pin tree. See the reply on the matrix-waiter thread.
…so both merge onto b10871
…nightly composes them from master
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 62cf502303
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| if (exact_pages && seq_id_src != seq_id_dst) { | ||
| LLAMA_LOG_ERROR("%s: exact concurrency does not support copying cells between " | ||
| "sequences (%d -> %d); ignoring the copy\n", | ||
| __func__, seq_id_src, seq_id_dst); | ||
| return; |
There was a problem hiding this comment.
On a supported CUDA setup with LLAMA_EXACT_CONCURRENCY=1 and -kvu, every cross-sequence llama_memory_seq_cp() returns here without populating the destination. This is not confined to the server's rejected n > 1 path: examples/parallel/parallel.cpp copies the shared system prompt at line 280, sets each client's n_past to that prompt's length at line 332, and then decodes only the suffix, so exact-mode outputs omit the system prompt while appearing successful. Allocate distinct destination pages or reject exact mode before callers enter flows that require sequence copying.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
The copy is not discarded silently: llama_kv_cache::seq_cp logs LLAMA_LOG_ERROR("exact concurrency does not support copying cells between sequences (%d -> %d); ignoring the copy") before returning, and llama_memory_hybrid::seq_cp does the same. llama_memory_seq_cp returns void, so an error log is the only signal the API allows, and aborting is what that path deliberately avoids. Rejecting before the caller enters the flow is already done for the only caller this repo ships: seq_cp is reached solely from server_slot::copy_state_to, whose only caller is the parent/child family loop, and a family only exists for n > 1, which exact mode turns into a 400 before the task is created. The upstream examples that copy a shared prefix predate the mode and I would rather not rewire four of them for an opt-in environment variable. What was missing is a regression test, so test-exact-pages now asserts that a refused copy leaves the destination empty and the source unchanged (#197 50aa4be, carried here as 1cfb42b), confirmed on a model that can run the mode.
| # Guards on the asynchronous per-sequence state transfer | ||
| # Skips itself on a backend that cannot copy asynchronously | ||
| llama_build_and_test(test-state-seq-copy.cpp LABEL "model" ARGS -m "${MODEL_DEST}") | ||
| set_tests_properties(test-state-seq-copy PROPERTIES FIXTURES_REQUIRED test-download-model) |
There was a problem hiding this comment.
This hunk registers five newly added files under tests/ (test-state-seq-copy.cpp, the three test-exact-* files, and test-server-tokens.cpp), but the repository requires maintainer approval before adding any new file in tests/*, and no such approval is documented in this change. Reuse existing test files for this coverage or obtain explicit maintainer approval before retaining the new files.
AGENTS.md reference: AGENTS.md:L84-L87
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
As the fork maintainer I approve these five standalone tests. test-exact-geometry is pure batch-shape arithmetic that links only common and needs no model or backend, so folding it into test-exact-pages would lose it on model-less and CPU-only CI, and test-backend-ops has no llama context at all; test-exact-buft needs ggml-backend plus the internal src/llama-impl.h and a device but no model, a third combination outside test-backend-ops; test-exact-pages needs a real context and a model and sits behind the download fixture under LABEL model; test-server-tokens links server-context and mtmd with the tools/server include path under LLAMA_BUILD_TOOLS, which no file in tests/ otherwise does; test-state-seq-copy drives llama_state_seq_copy_* on a real context, which test-state-restore-fragmented does not cover.
| // a host buffer type may quietly hand back ordinary memory when pinning is off, so believe the buffer that came back rather than the type | ||
| pinned = can_pin && ggml_backend_buffer_get_type(buf) == host_buft; |
There was a problem hiding this comment.
On a model split across multiple CUDA devices, this single host buffer is used by the per-device transfer backends, but the CUDA host buffer type allocates it with default cudaMallocHost semantics rather than portable pinning. Only the allocating CUDA context is guaranteed to regard that allocation as pinned; copies issued from the other devices can therefore stage or block even though this type comparison reports pinned = true, causing the server to enable --preempt-async while multi-GPU parks still stall decoding. Use portable pinned allocation, separate buffers per device, or verify pinning for every participating device.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
This is upstream ggml code the PR does not touch, and the flag is not needed. ggml_cuda_host_malloc calls cudaMallocHost, and the CUDA Runtime API's Unified Addressing section states that all host memory allocated through cudaMallocHost() and cudaHostAlloc() is always directly accessible from all devices that support unified addressing, regardless of whether cudaHostAllocPortable or cudaHostAllocMapped was specified, and that unified addressing is automatically enabled in 64-bit processes (https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__UNIFIED.html); the Programming Guide's Portable Memory section says the same, that the default benefits cover the allocating device and all devices sharing the unified address space. Every configuration the CUDA backend supports is a 64-bit process, so the buffer is portable everywhere the server runs and pinned = true is accurate. It is also the existing pattern: llama_context already takes the host buffer of the first device for intermediate state across a multi-device split, and the model loader picks the first device offering a host buffer type, so a real gap here would be an upstream defect in split mode generally.
This comment was marked as resolved.
This comment was marked as resolved.
Sorry, something went wrong.
…er exact concurrency
…est-exact-pages, the OpenVINO support probe's struct return, and park assertions that match the park log lines rather than the slot JSON (cherry picked from commit 025286e)
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 3e7d9530ac
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
|
|
||
| // DFlash drafting turns causal attention off on its draft context, which the paged attention needs; say so instead of asserting in the graph. DSpark is the same. | ||
| for (const auto type : params.speculative.types) { | ||
| if (type == COMMON_SPECULATIVE_TYPE_DRAFT_DFLASH || type == COMMON_SPECULATIVE_TYPE_DRAFT_DSPARK) { |
There was a problem hiding this comment.
With --spec-type draft-mtp on a shared-memory Gemma4 assistant, this guard admits a combination the exact KV allocator cannot execute normally: common_speculative_impl_draft_mtp::draft() intentionally submits every draft step at dp.n_past, while the paged allocator allows only one cell for a (sequence, position) pair. The second draft decode therefore returns a no-slot error on every iteration and truncates speculative decoding to one token. Reject exact mode after is_mem_shared is known, or support overwriting that draft cell.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
A shared-memory draft never claims a cell: apply_ubatch returns early when other is set and it is the only caller of exact_pages_claim, so every draft step sees the same free cell at n_past. The combination is also unreachable: ctx_other exists only for GEMMA4_ASSISTANT, EAGLE3 and DFLASH, the last two cannot produce an MTP context, and the Gemma 4 assistant is SWA, which exact mode refuses at cache construction.
…d-slot follower behind a slot that is seen busy, and give the notify resident time to decode on a loaded runner (cherry picked from commit 7529931)
(cherry picked from commit 8ec97d6)
… disconnect from raising test_a_resident_cycling_through_context_shifts_is_rotated_out_for_a_parked_head ran into the 600 s default request timeout on the ubuntu Server leg: the server was healthy and the third request came back with all 9000 tokens, the other two were still generating when the client gave up. The rotation is triggered by a 2 s wait, not by a token count, so the generation cannot be shortened without losing the rotation on a fast host; the requests get a longer timeout instead. test_a_resident_that_cannot_be_swapped_out_is_rotated_by_recompute stops the server with a request still in flight on purpose. Its thread raised the resulting ConnectionError as an unhandled thread exception, which is the noise that made the timeout above look like a crash. (cherry picked from commit 9bcd8ab)
…st's prompt stats A recompute restore re-enters prompt processing to put back the cells the park dropped. Those tokens are submitted with is_prompt set, so every one of them was added to n_prompt_processed and pushed t_prompt_last to the end of the replay, while n_gen deliberately carries across the park. The reported prompt length grew with every park, and the generation time covered only the tokens after the last re-prefill, so predicted_per_second was inflated by the ratio of the two. The replay is still counted in the server-wide prompt metrics, where it is real compute; it just no longer moves the slot's prompt count or the prompt/generation boundary. (cherry picked from commit d3833b0)
… streamed /v1/responses A non-streamed /v1/responses carries preempt in the response object, next to usage. The streamed one wrote it on the SSE data beside the response object, so a client that keeps the response of the response.completed event, which is the object the OpenAI SDK hands back, never saw it. It now sits in the same place either way, and unconditionally, as the non-streamed body already did. (cherry picked from commit e4168eb)
|
@codex security review |
This comment was marked as resolved.
This comment was marked as resolved.
Sorry, something went wrong.
…two requests overlapping (cherry picked from commit d14f07b)
|
@codex security review |
🛡️ Codex Security ReviewSecurity review completed. No security issues were found in this pull request. Reviewed commit: Only the user who started this review can view the report in Codex. ℹ️ About Codex security reviews in GitHubThis is an experimental Codex feature. Security reviews are triggered when:
Once complete, Codex will leave suggestions, or a comment if no findings are found. |
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: d8daa8fc51
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| return false; | ||
| } | ||
|
|
||
| if (llama_state_seq_get_data_ext(ctx_tgt, preempt_state_tgt.data(), size_tgt, id, LLAMA_STATE_SEQ_FLAGS_NONE) != size_tgt) { |
There was a problem hiding this comment.
On a CUDA server using --no-preempt-async or falling back to synchronous parking, a slot can be selected while processing an intermediate prompt batch, for which decode() deliberately skips llama_synchronize() because the batch has no output. These synchronous state calls then use CUDA copies on cudaStreamPerThread, while graph work runs on nonblocking streams, so copying the target or draft state can race the still-running KV writes; the corresponding synchronous restore at line 742 can similarly race preceding graph reads. Synchronize both contexts before the synchronous save/restore path, otherwise a parked request can resume from incomplete or corrupted KV state.
Useful? React with 👍 / 👎.
There was a problem hiding this comment.
Both synchronous calls already synchronize inside llama.cpp before any KV byte moves: llama_state_seq_get_data_ext and llama_state_seq_set_data_ext call ctx->synchronize() first (llama-context.cpp:4991 and :4996), which is ggml_backend_sched_synchronize over every backend, and on CUDA that is cudaStreamSynchronize on the main stream that the forked graph streams are joined back onto within graph_compute. The skipped synchronize after an output-free batch is recovered there, so a second llama_synchronize before the save or restore would be a duplicate round trip.
…refill (cherry picked from commit 5356710)
This comment was marked as resolved.
This comment was marked as resolved.
Sorry, something went wrong.
This comment was marked as resolved.
This comment was marked as resolved.
Sorry, something went wrong.
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: ae19539059
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| [](common_params & params, int value) { | ||
| params.preempt_ram_mib = value; |
There was a problem hiding this comment.
When --preempt-ram (or LLAMA_ARG_PREEMPT_RAM) is set to -2 or another unsupported negative value, this parser accepts it and preempt_ram_budget() treats every negative value as SIZE_MAX. A typo therefore silently removes the advertised RAM limit and can let parked states consume memory until allocation failure; reject values below the documented -1 sentinel.
Useful? React with 👍 / 👎.
| // [TAG_EXACT_CONCURRENCY] the page table of the paged path, which only the ordinary flash | ||
| // attention op carries: another op is free to keep a tensor of its own in the same slot |
There was a problem hiding this comment.
This newly added comment hard-wraps a single sentence between "flash" and "attention". Keep the sentence on one line or shorten it, as the repository explicitly prohibits hard-wrapping comments and splitting lines mid-sentence.
AGENTS.md reference: AGENTS.md:L77-L81
Useful? React with 👍 / 👎.
|
@codex review |
1 participant
The tree of #197 (its branch merged with the upstream release tag b10871) composed with the pinned commit of ggml-org#25731 (inkling, 946fc11), so the nightly's pin merge has one merge base after inkling.
History of this branch:
f32_pedanticguards inside the mul_mat dispatch server: preemption notices, asynchronous parks and exact concurrency together #197 restructured intoggml_cuda_mul_mat_path, and adds neighbours in llama-kv-cache.h and test-backend-ops.cpp at the same insertion points. The guard now lives inggml_cuda_mul_mat_path, so the switch, the invariant width and the column split pick the same implementation;ggml_cuda_mul_mat_id_needs_synckeeps both!f32_pedanticand server: preemption notices, asynchronous parks and exact concurrency together #197'sntokens; the flash attention precision call usesggml_prec_set_acc, which inkling introduces in place of the deprecated setter. CUDA build warning-free; test-exact-geometry, test-exact-buft, test-exact-pages pass; test-backend-ops MUL_MAT 1312/1312 and MUL_MAT_ID 930/930 on CUDA.Reproduced the preflight's resolve loop in a fresh clone of upstream on b10871 (temp replay of unsloth-pin-preflight.yml: fetch each pin,
merge --no-ffwith diff3, additive_merge.py on conflict): the thirteen current pins merge as the preflight reports, then this commit merges with two pure add/add hunks against earlier pins (the proc-address table entry next to ggml-org#24423's, the exact-concurrency refusals next to the kpool-dirty blocks in llama-memory-hybrid.cpp), which additive_merge.py resolves. merge_checks clean (23 python and 189 c++ files); pin_contract: all 14 pins intact, this one 25 files with new symbols and 6559 lines accounted for.The base branch upstream/b10871-inkling-25731 is the inkling pin pushed as a branch, so the diff here is #197's change set on top of inkling. Review happens on #197; this branch is only what gets pinned. Because the merge base is inkling's pinned commit, this pin has to be recomposed whenever ggml-org#25731 is repinned or leaves the set.
--preempt-ramnow defaults to 0, so a server launched without the flag behaves exactly as upstream, unified cache included: no planner, no pinned probe, no async copier, no startup notice, and the cache-full path fails every slot as upstream does (verified on GPU with four concurrent completions past a 512-cell pool: four context errors and zero preemption lines without the flag; four completions and 38 park and resume lines with--preempt-ram 1024). Exact concurrency stays an environment opt-in. Replay is clean on b10871 and on b10884, the newest aged tag at the time.src[5]only forGGML_OP_FLASH_ATTN_EXT, since inkling's banded op keeps its relative bias in the same slot (reverting the gate in the composed tree aborts on the banded MMA path;FLASH_ATTN_EXT_BANDED13/13 with it). The sequence-plane normalisation followsGGML_CUDA_BATCH_INVARIANTas its comment promises, not only the exact-mode variable. Replay clean on b10871 and b10884.