Repository navigation
[BREAKING][ops, model] refactor: replace OpSlot with VeomniOp - #1142
Merged
Coach257 merged 88 commits intoOct 18, 2026
Merged
Conversation
|
Important Review skippedToo many files! This PR contains 747 files, which is 447 over the limit of 300. To get a review, reduce the PR to 300 files or fewer by splitting it into smaller PRs or changing its base branch. Usage-priced reviews support at most 300 files. ⚙️ Run configuration
⛔ Files ignored due to path filters: 64
📒 Files selected for processing: 747
You can disable this status message by setting the
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
Collaborator
Coach257
reviewed
Sep 3, 2026
Coach257
reviewed
Sep 3, 2026
Coach257
reviewed
Sep 3, 2026
Coach257
reviewed
Sep 3, 2026
Coach257
reviewed
Sep 7, 2026
…omniOp Authors register raw forward/backward or an opaque wrapper. The registry generates the autograd Function, keys rows by (op, variant, impl, device), and hands modeling a local VeomniOp. Delete the old kernel_registry, dispatch, and config singleton so the public API has one registration path.
…l loss Move the common tensor ops onto OP_REGISTRY with eager and vendor rows. Keep CE and load-balancing at token level; shifting, SP reduction, and log-prob policy stay out of the kernel rows. Include LayerNorm so async Ulysses does not depend on Apex.
Add the remaining compute families as registry rows, including LoRA expert variants and device-specific implementations. Compound or opaque ops keep raw nested calls where a clean forward/backward split exists.
Install stock and veomni_* attention into the Transformers registries. Keep mask builders as a separate shape/native API. Async Ulysses QKV/O uses standard and dit variants that call nested raw pairs.
…andles Model construction, checkpoint convert/load, and loss helpers resolve ops from config and store VeomniOp on the instance. There is no process-wide bind. Shared tests cover registry build, config scope, and handle isolation.
Point every supported text, VL, and Omni family at instance-local VeomniOp handles. Keep patchgen configs and generated modeling in the same commit so the registered classes match the new call sites.
Consume Wan, Flux, Qwen Image, LTX, MiniMax, and MoVQGAN through the same VeomniOp handles as the transformer families, including DiT async Ulysses and family-specific attention wiring.
Update trainer, sequence-parallel, LoRA, data, and YAML defaults to the new op names and instance-local handles. Drop leftover slot/facade rewrites so callers cannot select interned names that no longer exist.
Describe OpEntry, VeomniOp, and selection in current docs and skills. Point patchgen, device-api, and unit-test workflows at the canonical ops and models packages.
Zx55
force-pushed
the
czr/kernel-registry-refactor
branch
from
September 20, 2026 03:51
520a68c to
c29b58f
Compare
Zx55
marked this pull request as ready for review
September 20, 2026 04:11
… and NPU checks Co-authored-by: Cursor <cursoragent@cursor.com>
… checks Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
Pin the SwiGLU routing scale to the activation dtype so batch-invariant linear does not mix fp32 router scores with bf16 expert weights. Align the Wan padded Ulysses reference to the same mask, skip Qwen-Image isolation on NPU CPU tensors, and give DeepSeek-V3 e2e the same bf16 mixed-precision extra_args as V4. Co-authored-by: Cursor <cursoragent@cursor.com>
… on compute dtype Drop the GPU eager-MoE override so DeepSeek-V3 e2e stays on fused_triton with the original EP=2 matrix. Cast fp32 router scores onto the activation dtype at the experts call site and in the Triton wrapper so group-gemm does not see a dtype split. Co-authored-by: Cursor <cursoragent@cursor.com>
Coach257
reviewed
Oct 17, 2026
Co-authored-by: Cursor <cursoragent@cursor.com>
Co-authored-by: Cursor <cursoragent@cursor.com>
…s, and NPU detection Co-authored-by: Cursor <cursoragent@cursor.com>
…h VeomniOp Co-authored-by: Cursor <cursoragent@cursor.com>
Coach257
approved these changes
Oct 18, 2026
Estrellajer
added a commit
to Estrellajer/VeOmni
that referenced
this pull request
Oct 19, 2026
- Implement shared-prefix training for Qwen3.5 and Qwen3.5 MoE on GPU and NPU - Adapt GatedDeltaNet and Attention layers to VeomniOp interface (ByteDance-Seed#1142) - Fork recurrent state for GatedDeltaNet at prefix boundary - Gather key/value states for Attention across prefix and suffix - Add shared prefix plan layout and compaction utilities - Add comprehensive unit tests and documentation
wangyier7
added a commit
to wangyier7/VeOmni
that referenced
this pull request
Oct 19, 2026
…tern Follow-up to the origin/main merge (ByteDance-Seed#1142 OpSlot -> VeomniOp refactor): - KimiRMSNorm / SituAndMul: __init__-bound instance-local VeomniOp handles (rms_norm/standard, situ_and_mul/kimi_k3) replacing the OpSlot guards - PatchedKimiK3Experts: branchless dispatch through an instance-local moe_experts VeomniOp (kimi_k3 variant for SituGLU, standard otherwise); the eager expert loop is now a registered row, not an in-rule fallback - CE loss: ForCausalLMLoss bound to a cross_entropy_loss/standard VeomniOp via functools.partial on both KimiLinearForCausalLM and KimiK3ForConditionalGeneration (variant renamed causal -> standard) - chunk_kda: name hijack now dispatches through a lazily-resolved chunk_kda/kimi_k3 VeomniOp (first-call resolution follows the ops config installed at model-build time) - NPU config re-registers the new __init__ overrides; regen both targets
5 tasks done
Coach257
added a commit
that referenced
this pull request
Oct 20, 2026
…le via VeomniOp Merge #1314 (including #1142 VeomniOp refactor) into omniv2 alongside the per-module SP/CP/CP*SP support, and migrate SeedOmni modules onto explicit VeomniOp selection from self.config.ops_implementation: - BagelQwen2MoTAccelerated, BagelSiglipNavit, TextEncoderAccelerated, JanusVqvaeAccelerated, Qwen3MoeLlm: replace module-level OpSlot / global LOSS_MAPPING / direct kernel dispatch with instance-local VeomniOp handles constructed directly from config.ops_implementation.get(...). - JanusSiglipConfig: declare sub_configs = {"vision_config": JanusVisionConfig} so _attn_implementation propagates to the nested vision config. - OmniModel._load_modules, ModuleRuntime._init_eager_inference: drop global set_ops_config / bind_ops_to_modeling calls.
5 tasks done
Open
4 tasks done
yeyeyeping
added a commit
to yeyeyeping/VeOmni
that referenced
this pull request
Oct 21, 2026
…MiniMax M3 VL support MiniMax M3 VL (MiniMaxAI/MiniMax-M3) support rebuilt on main after the OpSlot -> VeomniOp refactor (ByteDance-Seed#1142), the moe_experts reorganisation, and the test-tree reorganisation. - Modeling: patchgen GPU/NPU configs construct instance-local VeomniOps: rms_norm/offset (fp32 1 + weight), cross_entropy_loss via ForCausalLMLoss, moe_experts/swiglu_oai, minimax_sparse_attention indexer/attention, and attention/standard for the vision tower. Packed-only language attention, Ulysses with per-KV-group index-query exchange, frozen indexer, vision attention guard for varlen and veomni_* backends, router aux loss disabled. - Ops: minimax_sparse_attention eager rows (npu_fusion_attention on Ascend); moe_experts/swiglu_oai rows (eager, fused_triton, fused_npu) reusing the standard merged-fc1 Triton/NPU/EP paths through an optional swiglu_alpha; SwiGLU-OAI helpers in moe_experts/shared/swiglu.py; Ulysses gather_seq_scatter_kv_heads. - Data/trainer: MiniMax processor, chat template and data transform; shared VLM media loading; vision_tower lookup and vit_lr grouping. - Checkpoint: runtime converter for the public checkpoint layout. - Tests, configs and docs for the above.
yeyeyeping
added a commit
to yeyeyeping/VeOmni
that referenced
this pull request
Oct 22, 2026
…MiniMax M3 VL support MiniMax M3 VL (MiniMaxAI/MiniMax-M3) support rebuilt on main after the OpSlot -> VeomniOp refactor (ByteDance-Seed#1142), the moe_experts reorganisation, and the test-tree reorganisation. - Modeling: patchgen GPU/NPU configs construct instance-local VeomniOps: rms_norm/offset (fp32 1 + weight), cross_entropy_loss via ForCausalLMLoss, moe_experts/swiglu_oai, minimax_sparse_attention indexer/attention, and attention/standard for the vision tower. Packed-only language attention, Ulysses with per-KV-group index-query exchange, frozen indexer, vision attention guard for varlen and veomni_* backends, router aux loss disabled. - Ops: minimax_sparse_attention eager rows (npu_fusion_attention on Ascend); moe_experts/swiglu_oai rows (eager, fused_triton, fused_npu) reusing the standard merged-fc1 Triton/NPU/EP paths through an optional swiglu_alpha; SwiGLU-OAI helpers in moe_experts/shared/swiglu.py; Ulysses gather_seq_scatter_kv_heads. - Data/trainer: MiniMax processor, chat template and data transform; shared VLM media loading; vision_tower lookup and vit_lr grouping. - Checkpoint: runtime converter for the public checkpoint layout. - Tests, configs and docs for the above.
This was referenced Oct 23, 2026
sc1915
pushed a commit
to sc1915/VeOmni
that referenced
this pull request
Oct 23, 2026
…nce-Seed#1142) Co-authored-by: Cursor <cursoragent@cursor.com>
5 tasks done
5 tasks done
38 tasks
6 of 8 tasks
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Merge stack
Ops PRs that land before this one, because they change kernels this refactor moves or add kernels on the old
OpSlotAPI. Every other open ops PR rebases ontoVeomniOpafter this merges, and carries a note pointing here.Already on
main— port into this branch on the next rebase:b199a26:cannrows fordsa_indexer/dsa_attention/mhc, and the DeepSeek-V4 NPU patch_gen dispatches throughVeomniOp.53d5383ontogated_delta_rule/vendor/; its regression test moves totests/ops/gated_delta_rule/.53d5383: a newqsa_attentionop witheagerandtilelangrows, and the qwen4_exp patch_gen keepsmain's Ulysses logic on instance-localVeomniOps.triton-ascendin the NPU unit and e2e workflows. Ported inb60b81176: installtriton==3.2.0, thentriton-ascend==3.2.1, and checktriton._C.libtriton.ascend.b60b81176onto the registry paths.torch_npurms_norm_gatedusesnpu_swiglufor SiLU and the eager fp32 gate multiply for sigmoid. Tritoncausal_conv1dandchunk_gated_delta_rulemove a CPUcu_seqlensonto the activation device. The backward autotune key is['H', 'K', 'V', 'BT']. The new tests callresolve_op. NPU training defaults and the Qwen4-Exp e2e smoke usetorch_npu/triton, with SP=2.4b21b4c6d.set_ulysses_sequence_parallel_groupis gone. SP and CP tests build real states with_init_parallel_stateand enter them withuse_parallel_state.4140f586contovendor/triton/.NT,N,B, andTOTAL_TASKSleavetl.constexprand joindo_not_specializeinchunk_delta_h,chunk_o,chunk_scaled_dot_kkt, andwy_fast.causal_conv1dbackward over-countingdh0when the tile is narrower than the conv window. Merged in4140f586c.mask_head_rowsand thedy0loads are bounded byT_len. Withinitial_state,BTis at leastnext_power_of_2(W), andW > 8raises. The packed short-segment regression is intests/ops/gated_delta_rule/test_gated_delta_rule.pyand calls the vendored impl, because the registry wrapper still passesinitial_state=None.tests/ops/test_npu_kernels.pystays deleted.Deliberately not in this stack, rebasing onto
VeomniOpafterwards instead: #1223 (the sigmoid-vs-npu_swiglusplit is already ported from #1309; the remaining device/dtype and invalid-gate contracts edit the deletednpu_rms_norm_gated.py), #966 (910B4-1BV=64profile — the root cause has no hardware evidence yet, and it edits files this PR deletes), #1306 (merged-MoE backward activation lifetime — a memory optimization, not a correctness fix), #1199 (expert-parallel numerics — too large to sit in front of the refactor; it changes routing-weight arithmetic for every Triton MoE user), #1200 (compiled-training precision — two of its three code changes are already resolved here, see the note on that PR).Why
VeOmni previously selected optimized operations through several independent mechanisms, including
OpSlot, kernel factories, global function replacement, loss mapping mutation, and attention-specific registries.This made implementation selection difficult to reason about and prevented different model instances from safely selecting different implementations.
This PR consolidates tensor-level operations under
veomni.ops. Models hold instance-localVeomniOphandles. Model-specific argument processing and loss policy remain inveomni.models.What does this PR do?
The refactor creates one path from implementation registration to instance-local model execution:
flowchart TB subgraph Registration["1. Operation registration"] direction LR Raw["Raw forward/backward pair"] Opaque["Opaque eager or vendor wrapper"] Register["register_op(...)"] Entry["OpEntry<br/>wrapper + optional raw pair"] Registry["OP_REGISTRY<br/>(op, variant, implementation, device)"] Raw --> Register Opaque --> Register Register --> Entry Entry --> Registry end subgraph Runtime["2. Model construction and ordinary execution"] direction LR YAML["model.ops_implementation"] Config["OpsImplementationConfig"] Model["Model __init__"] Handle["VeomniOp(op, variant, impl)"] Resolve["Resolve row and<br/>check hardware"] Call["entry.wrapper(...)"] Kernel["Selected implementation<br/>PyTorch / Liger / Triton / NPU / MLU"] YAML --> Config Config --> Model Model --> Handle Handle --> Resolve Registry --> Resolve Resolve --> Call Call --> Kernel end subgraph Compound["3. Compound operation execution"] direction LR Outer["Outer raw forward/backward"] Inner["resolve_inner_op(...)"] RawCall["Call inner raw forward/backward<br/>without nesting an autograd wrapper"] State["Flatten and restore SavedState"] Outer --> Inner Registry --> Inner Inner --> RawCall RawCall --> State endThe important properties are:
OP_REGISTRY, including an expliciteagerrow.VeomniOp; constructing one model no longer changes another model's implementation.Function.apply.veomni.ops; model-specific input adaptation and loss policy live inveomni.models.Process-wide integration is limited to installation work such as registering VeOmni attention names with Transformers. Attention masks, batch-invariant patches, and model-side loss policy are not registered operation rows.
HF-interface families bind attention from
config._attn_implementation. Families that are not on that path (Wan, MiniMax, Flux, Qwen Image, LTX) define no model-localeager_attention_forward, so they bind throughresolve_op_impl("attn_implementation", eager_as="veomni_sdpa").The complete operator catalog and extension guide are documented in
veomni/ops/README.md.Authoring, registration, and modeling usage of
VeomniOp/OP_REGISTRYare documented inveomni/ops/README.mdanddocs/design/op_selection.md.Config migration
model.ops_implementation.*keeps its field names. Values renamed sincemainstill parse for one transition window.OpsImplementationConfigmaps them to the current name and logs a deprecation warning that names the replacement. The table is also indocs/usage/arguments.md.mainattn_implementationveomni_*_with_sp(FA2/FA3/FA4, Hub forms, Flex, Magi)_with_sp. Ulysses now lives inside the adapter.moe_implementationfusedmainresolved it to on that host:fused_quack(GPU),fused_npu(NPU),fused_mlu(MLU with Apex) orfused_triton(MLU without Apex).moe_implementationfused_mlu_tritonfused_tritoncross_entropy_loss_implementationnpuchunk_loss. Themainentry was an alias ofchunk_loss, not a third-party NPU kernel.rms_norm_implementation,rotary_pos_emb_implementation,rotary_pos_emb_vision_implementation,rms_norm_gated_implementationnputorch_npucausal_conv1d_implementation,chunk_gated_delta_rule_implementationnputriton, the vendored Triton-Ascend kernelchunk_gated_delta_rule_implementationnpu_ascendcfla_npurms_norm_implementationis shared by every RMSNorm variant, so one value is also mapped per variant.unweighted(Flux, LTX-2.3, DeepSeek-V4) has no Ascend kernel, sotorch_npuand the oldnpuboth resolve toeagerfor that variant, because themainrow of that name was PyTorch math. Consumers passresolve_op_impl(field, variant=...), and the mapping is logged once.rms_norm_implementation='qwen3_5'is rejected at parse time.qwen3_5was the RMSNorm variant name (nowoffset), not an implementation value. Other unknown values fail at model build, and the error lists the registered implementations.Host defaults
The seven general-purpose fields (
moe,cross_entropy_loss,rms_norm,swiglu_mlp,rotary_pos_emb,rotary_pos_emb_vision,load_balancing_loss) now default toNone, andOpsImplementationConfigfills an unset field with the host default. The filled values are the same asmainon GPU, NPU and MLU. The table is indocs/usage/arguments.md.maintreated a value equal to the GPU dataclass default as unset and rewrote it on NPU and MLU, even when the YAML set it explicitly. Explicit values are now never rewritten:liger_kernel,fused_triton, orload_balancing_loss_implementation: triton) raises at config validation instead of silently selecting the NPU kernel. GPU YAMLs that pin these keys must drop or replace them for NPU runsrotary_pos_emb_vision_implementation: eagerstayseager.mainrewrote it tonpuliger_kernelis kept and fails at model build, wheremainrewrote it toeager. An explicitfused_tritonMoE staysfused_tritoneven with Apex, wheremainswitched it tofused_mluOperator surface differences from
mainExcept for the differences below, the refactored operations are intended to preserve the tensor contracts, numerical behavior, gradients, and hardware support of
main. Moving an existing implementation into the new registry is not considered a behavior change.mainunweightedrms_norm_implementation='torch_npu'(or the oldnpu) resolves toeagerfor this variant, because the previous NPU-labelled path was PyTorch math and did not call a distincttorch_npukernel.deepseek_v4deepseek_v4/{eager,liger_kernel,torch_npu}.mainrouted V4 throughstandard. Thetorch_npurow calls the sametorch_npu.npu_rms_normkernel thatmainreached throughstandard.offsetqwen3_5variant is renamed tooffset. Rows areoffset/{eager,liger_kernel,torch_npu}. Consumers are Qwen3.5, Qwen3.5-MoE, Gemma3, and Qwen4-Exp.group_sizegroup_size. Fused weighted rows delegate grouped calls to eager. Only Qwen4-Exp passes a non-Nonevalue (hc_norm,norm_key,norm_query,norm_conv).main's local reshape; other callers omit it.standard/{eager,chunk_loss,liger_kernel}. The previousnpuentry was an alias ofchunk_loss.npunow parses as a deprecated config alias ofchunk_losson any host. Causal and sequence-classification policy remains model-side.rope_visionfamily. Rank-3 vision and text full RoPE both userope/full.rotary_pos_emb_vision_implementationremains temporarily as an independent configuration choice.rms_norm_gated/standard/{eager,fla,torch_npu}acceptactivation="sigmoid", which Qwen4-Exp selects throughoutput_gate_type.torch_npufuses SiLU withnpu_swigluand uses the eager fp32 gate multiply for sigmoid.maindispatchesqsa_attention_implementationinside the modelingeager_attention_forward. Here it is theqsa_attention/standard/{eager,tilelang}op, which takes the compact global-index selection. The indexer and attention still compute the selection and run the Ulysses exchange in modeling.tilelangraises on an attention mask and falls back toeagerwith a warning for nonzero dropout, as onmain. The non-QSA eager fallback ineager_attention_forwardis gone because the vision tower binds its ownattentionop.rope/partial, so the GPU defaultliger_kernelapplies wheremainused the HF eager function.activation="silu", matching the model-path normal operation. The previous default wasNone.main'snpurows are renamed after what they call:torch_npu(RMSNorm, RoPE, gated RMSNorm) andtriton(causal Conv1D, chunk gated delta rule).npu_ascendcbecomesfla_npu.main's NPU patch_gen calls the CANN Lightning Indexer, sparse FlashMLA and mHC kernels inside modeling code. Here they arecannrows ofdsa_indexer,dsa_attentionandmhcwith the same structured contract as the GPU rows. The attention row converts global top-k indices to the segment-local indices CANN TND expects.mhc/head/cannis the eager math, as onmain, since CANN has no head kernel.main's unreferencednpu_sparse_flash_mla_with_indexer_loss.pyis not ported.main.attention_maskis still ignored, and calls without compressed rows still fall back to eager. The adapters no longer pre-check dtype or shape limits (for example the mHC bf16 check) and leave them to the kernel.fused_triton,fused_quack,fused_npu, andfused_mluvalues exposed by configuration._with_sp:veomni_flash_attention_*,veomni_flex_attention,veomni_magi_attention.sdpais rewritten toveomni_sdpa. Packed/varlen metadata is dropped when it describes at most one segment per row or a 3-D/4-D mask isolates the segments.packed_causal_mask(andsliding_window_maskwithcu_seqlens) builds a dense block-diagonal mask for SDPA. EFFICIENT/MATH backends are pinned only for masked calls.mainignored the metadata and attended across samples. Masked forward matches HFsdpabitwise. Backward is not bitwise where HF's default dispatch picks another backend (cuDNN on H20, up to 3.9e-3 on tiny bf16 models).eagerin non-HF familieseagertoveomni_sdpa, and useveomni_sdpaas the masked fallback.main's eager path.Validation
The ops suite covers registry contracts, public signatures, forward and backward parity, hardware requirements, and optional dependency handling. The models suite covers construction, instance-local selection, family parity, and model-facing operator integration.
Trainer composition,
ModelArgumentspath fallback, and SPReduceLosslive with their owners rather thantests/models.Latest local GPU result on 8× H20 for
pytest tests/ops tests/models tests/lora tests/arguments, before the four fp32 loader cases were added to the oracle:make check-patchgenand the pre-commit hooks pass.For the
53d5383merge of #1197 / #970 / #1293 / #1300, the local 8× H20 results are as follows. The NPU paths of #970 and the Qwen4-Exp NPU generated file wait for CI.pytest tests/ops tests/arguments tests/models tests/special_sanity tests/tools, plus the SeedOmni HF-layout and FLOPs tests, gives1723 passed, 454 skippedtests/ops/qsa_attentionandtests/ops/gated_delta_rulegive48 passed, with NPU/MLU-only skipstests/parallel/ulysses/test_qwen4_exp_ulysses.pyon 4 GPUs gives54 passed, including SP=4test_qwen4_exp_training_smoke(SP=2, EP=2, PLE=2,qsa_attention_implementation=tilelang) passes51d3898merges #1279, the SeedOmni offline-cache workflow, which changes no op code. Only import paths andhf_layout.mdanchors needed resolving.tests/seed_omniandtests/checkpoints/test_{checkpoint_callback,dcp_checkpointer}.pygive2260 passed, 456 skipped, 4 failedtest_dcp_checkpointer.pycases that calltorch.cpu.empty_cache. They fail the same way onmainwith GPUs hiddenb60b81176merges #1307 / #1309.pytest tests/ops/gated_delta_rule/test_gated_delta_rule.py tests/arguments/test_ops_implementation_aliases.py, with the NPU and FLA cases deselected, gives31 passed, 3 skipped. The skips are the new NPU sigmoid and CPU-cu_seqlenscasesmake commitpassede94f4950esizes the NPUcausal_conv1dparity case to the kernel tile (D=256). The forward withoutinitial_statetiles the channel axis at 256. That case skips locally because this host has no NPU.4b21b4c6dmerges #1311.make commitpassedruff. GPU parity was not re-run4140f586cmerges #1264 / #1265.make commitpassedUsage
Models construct an operation handle once and call it directly:
The exact registered and currently available implementations can be inspected through the registry: