Repository navigation
Conversation
…ip bitwise) Port of miles@ba28991834 tests/test_hf_to_megatron.py conventions. Adds tests/converters/harness.py with assert_roundtrip_bitwise and assert_hf_logits_preserved plus tiny 2-layer configs; smoke coverage on qwen3. Day-0 converter PRs (DSv4/v4.1, GLM5-Next, Kimi-K3, Qwen3.8-Next) plug in here. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/mbridge/{__init__,linear_attn,patch/}
and the _load_hf_weights_with_mbridge path in miles/backends/megatron_utils/checkpoint.py.
New --hf-to-megatron-load-mode {builtin,mbridge} (default builtin); mbridge pinned
to 89eb1088 in Dockerfiles with --no-deps, recorded in requirements.txt.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles/kernels/attention/dsa/{__init__,indexer,topk,triton/}.
TileLang backends skipped (default score_backend=triton; torch recompute backward
fallback); sglang canonical topk falls back to torch with warning; sglang envs
replaced by VIME_DSA_TOPK_* env vars.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles/kernels/attention/dsa/{kpool,sparse_attention}.py.
kpool is pure Triton; sparse_attention keeps guarded flash_mla/cudnn backends and
raises a clear ImportError for the unported tilelang backends. GPU tests adapted:
no miles ci_register, tilelang parametrizations gated on availability.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles/utils/replay_base.py. Engine-agnostic BaseReplayManager + IndexerReplayManager; RoutingReplayManager intentionally not ported (vime/utils/routing_replay.py is ahead). current_device via vime.utils.accelerator. Needed by the DSv4 indexer (v4_indexer.py). Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/models/deepseek_v4/{arguments.py,
ops/{rope,norm,precision_aligned_ops,utils,qat}.py}.
qat.py uses a pure-torch UE8M0 fp8-simulate fallback (tile_kernels / A2b quant
kernels not required); precision_aligned_ops gains a CPU upcast path.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/models/deepseek_v4/ops/{thd_utils,cp_utils,cp_row_balance}.py
with their CPU test suites (tests/fast/test_dsv4_thd.py, tests/fast/.../test_cp_row_balance.py,
plus tests/fast/dist_utils.py gloo helpers).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/models/deepseek_v4/ops/{v4_indexer,compressor}.py
plus the torchrun CP-balance test (tests/fast-gpu/test_dsv4_indexer_cp_balance.py,
GPU-only, NUM_GPUS=2).
topk backend reads config.miles_dsa_topk_backend with a 'canonical' default
(vime has no typed-config flag; canonical falls back to torch on vLLM).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/models/deepseek_v4/deepseek_v4.py. miles_dsa_topk_backend reads via getattr with miles' 'canonical' default (vime has no typed-config flag). Spec wiring tested with mocked megatron spec registry. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/mbridge/deepseek_v4.py and miles/backends/megatron_utils/megatron_to_hf/deepseekv4.py. New builtin hf_to_megatron loader (inverse of the exporter, incl. packed hyper-connection alpha split); megatron_to_hf dispatch gains 'deepseekv4' before the deepseekv3 family; bridge registered as model_type deepseek_v4. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 the return_indexer_topk contract (miles/rollout/session/config.py,
miles/rollout/generate_utils/generate_endpoint_utils.py) onto vLLM.
New docker/patch/latest/vllm-indexer-topk.patch: --enable-return-indexer-topk flag,
IndexerTopkCapturer hooked into SparseAttnIndexer (per-layer slots, spec-rejected
suffix dropped), per-request chunks on EngineCoreOutput.indexer_topk, base64 .npy
on the token_in_token_out choice. vime side: --use-rollout-indexer-replay +
--rollout-indexer-topk-num-streams flags, decode in vllm_rollout/vllm_streaming_rollout,
Sample.rollout_indexer_topk with shape/continuation validation.
Patch verification (fresh worktree of vllm@43b4aaea3 + all docker/patch/latest
patches in Dockerfile order, then):
git apply --check -v docker/patch/latest/vllm-indexer-topk.patch
result: applies cleanly.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Patch regenerated at its Dockerfile apply position (after vllm-draft-weights-cpu-backup);
indexer_topk is now the true last EngineCoreOutput field. Verified on a fresh worktree of
vllm@43b4aaea3 with the exact loop:
git apply vllm-rust-{frontend,observability,multimodal}.patch
git apply --include='rust/*' vllm-rust-server-info.patch
git apply -v vllm-pull_weights.patch && git apply -v --allow-empty vllm.patch
git apply -v <vllm-score-centering ... vllm-draft-weights-cpu-backup, Dockerfile order>
git apply --check -v docker/patch/latest/vllm-indexer-topk.patch
-> all hunks "Checking patch ..." OK, exit code 0
ROCm: intentionally not applied in Dockerfile.rocm (comment added there).
Port of miles@ba28991834 miles/backends/training_utils/replay/{routing_replay,data}.py
(indexer subset) onto vime's routing-replay structure.
New vime/backends/megatron_utils/indexer_replay.py: fill_indexer_replay +
prepare_indexer_topk_for_replay (THD token-position rebase, zigzag CP, SP slice,
-1 padding) + register_replay_list_sequential. Model build enables
IndexerReplayManager via VIME_ENABLE_INDEXER_REPLAY (ray actor_group env);
actor stage transitions mirror ROUTING_REPLAY_STAGE.
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/models/{linear_attn.py,cp_utils.py},
miles/backends/megatron_utils/{fp32_param_utils.py,megatron_to_hf/linear_attn_layout.py}.
linear_attn.py keeps vime's existing hf_attention.py untouched; forced-dtype attr
renamed to _vime_forced_param_dtype and enforce_marked_param_dtypes wired into
setup_model_and_optimizer (no-op when nothing is marked). Layout tests ported
verbatim from miles (tests/fast/models/test_linear_attn_layout.py).
Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Port of miles@ba28991834 miles_plugins/models/kimi_k3/, miles_plugins/mbridge/kimi_k3.py, miles/backends/megatron_utils/megatron_to_hf/kimi_k3.py. New builtin hf_to_megatron/kimi_k3.py loader (inverse of the exporter, nested language_model.model.* HF names, KDA/MLA layer split); megatron_to_hf dispatch gains 'kimik3' ahead of the kimi->deepseekv3 family; bridge registered as model_type kimi_k3. KDA vs MLA g_proj collision avoided by layer-type split. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
There was a problem hiding this comment.
Code Review
This pull request introduces support for Day-0 models (DeepSeek-V4 and Kimi-K3) by adding custom model plugins, Triton-based DSA (DeepSeek Sparse Attention) kernels, and an opt-in mbridge loading mode for HuggingFace checkpoints. It also implements a rollout indexer top-k replay mechanism to record and replay indexer decisions during training. The review feedback highlights three critical issues: an incorrect offset calculation in the indexer replay preparation when context parallelism is enabled, an erroneous tensor-parallel sharding of replay data under sequence parallelism, and a missing mask for invalid columns in the fallback PyTorch top-k implementation when sglang is unavailable.
| # Rebase per-sample KV positions onto the packed training sequence. | ||
| local = torch.where(local != -1, local + offset, local) | ||
| offset += local.shape[0] | ||
| sliced.append(local) |
There was a problem hiding this comment.
When context parallelism is enabled (cp_size > 1), local is a sliced version of the sample's top-k indices. Incrementing offset by local.shape[0] (the local sliced size) rather than token_len (the global size of the sample) will cause incorrect indexing for subsequent samples in the packed sequence, as their keys in the all-gathered k are located after the entire previous sample. We should increment offset by token_len instead.
| # Rebase per-sample KV positions onto the packed training sequence. | |
| local = torch.where(local != -1, local + offset, local) | |
| offset += local.shape[0] | |
| sliced.append(local) | |
| # Rebase per-sample KV positions onto the packed training sequence. | |
| local = torch.where(local != -1, local + offset, local) | |
| offset += token_len | |
| sliced.append(local) |
| if sequence_parallel and tp_size > 1: | ||
| seqlen = replay_data.size(0) | ||
| assert seqlen % tp_size == 0 | ||
| replay_data = replay_data[seqlen // tp_size * tp_rank : seqlen // tp_size * (tp_rank + 1)] |
There was a problem hiding this comment.
The DSA indexer always gathers the inputs to the full sequence length when sequence_parallel is enabled, meaning the attention computation and the indexer run on the full sequence length on each TP rank. Sharding replay_data by TP rank here will cause a shape mismatch and assertion failure during training. This sharding block should be removed.
return replay_data| def torch_dsa_topk(logits: torch.Tensor, topk: int, row_starts=None, row_ends=None) -> torch.Tensor: | ||
| # torch.topk errors on k > n; a context shorter than index_topk is legal, so pad with -1 | ||
| k = min(topk, logits.shape[-1]) | ||
| score, indices = torch.topk(logits, k, dim=-1) | ||
| indices = indices.to(torch.int32) | ||
| indices = indices.masked_fill(score == -torch.inf, -1) | ||
| if k < topk: | ||
| pad = indices.new_full((*indices.shape[:-1], topk - k), -1) | ||
| indices = torch.cat([indices, pad], dim=-1) | ||
| return indices |
There was a problem hiding this comment.
When sglang is not installed and topk_backend == "canonical", logits contains garbage values in the invalid columns because clean_logits is set to False. Since torch_dsa_topk does not use row_starts and row_ends to mask out invalid columns, it will incorrectly select garbage indices. We should mask out the invalid columns using row_starts and row_ends if they are provided.
| def torch_dsa_topk(logits: torch.Tensor, topk: int, row_starts=None, row_ends=None) -> torch.Tensor: | |
| # torch.topk errors on k > n; a context shorter than index_topk is legal, so pad with -1 | |
| k = min(topk, logits.shape[-1]) | |
| score, indices = torch.topk(logits, k, dim=-1) | |
| indices = indices.to(torch.int32) | |
| indices = indices.masked_fill(score == -torch.inf, -1) | |
| if k < topk: | |
| pad = indices.new_full((*indices.shape[:-1], topk - k), -1) | |
| indices = torch.cat([indices, pad], dim=-1) | |
| return indices | |
| def torch_dsa_topk(logits: torch.Tensor, topk: int, row_starts=None, row_ends=None) -> torch.Tensor: | |
| if row_starts is not None and row_ends is not None: | |
| columns = torch.arange(logits.shape[-1], device=logits.device) | |
| in_range = (columns >= row_starts.long().unsqueeze(-1)) & (columns < row_ends.long().unsqueeze(-1)) | |
| logits = logits.masked_fill(~in_range, float("-inf")) | |
| # torch.topk errors on k > n; a context shorter than index_topk is legal, so pad with -1 | |
| k = min(topk, logits.shape[-1]) | |
| score, indices = torch.topk(logits, k, dim=-1) | |
| indices = indices.to(torch.int32) | |
| indices = indices.masked_fill(score == -torch.inf, -1) | |
| if k < topk: | |
| pad = indices.new_full((*indices.shape[:-1], topk - k), -1) | |
| indices = torch.cat([indices, pad], dim=-1) | |
| return indices |
Documentation build overview
62 files changed ·
|
Ports miles' day-0 model support: opt-in mbridge HF→Megatron load, DSA indexer/topk/kpool/sparse-attention Triton kernels, the DeepSeek-V4 plugin (ops, THD packing, CP row-balance, indexer + compressor, attention, converter, bridge), DSA indexer top-k replay (vLLM channel via
vllm-indexer-topk.patch+ trainer wiring), and the Kimi-K3 plugin (line-for-line with miles).Source:
radixark/miles@ba28991834. Port is additive (new modules + newadd_*_arguments()groups); slime-origin files are not restructured so future[Sync] Mirror SlimePRs stay clean. One commit per sub-PR; review each commit separately.Commits:
Verification: F07-0 converter harness (2-layer logits allclose + round-trip bitwise). The vLLM patch applies in full Dockerfile order (after
vllm.patch); docker build still pending. Follow-ups: DSv4.1, GLM5-next, Qwen3.8-next, ROCm side of the indexer patch.Merge order of the miles-port series: F10 → F02 → F06 → F11a → F11b → F05 → F09 → F04 → F03 → F07 → F08 → F01.
train.py/vime/utils/arguments.pyare touched by several branches; later ones are rebased after each merge.🤖 Generated with Claude Code