feat(enflame): port GCU300 enablement bundle to main - #432
Open
tengqm wants to merge 2 commits into
Open
Conversation
![github-advanced-security[bot]](https://avatars.githubusercontent.com/in/57789?s=60&v=4){kind=small}
| import vllm.v1.worker.block_table as bt | ||
|
|
||
| bt.BlockTable.compute_slot_mapping = _compute_slot_mapping_torch | ||
| _patched = True |
| apply_chunk_delta_h_gcu_patch() | ||
| apply_fused_recurrent_packed_decode_gcu_patch() | ||
| apply_slot_mapping_gcu_patch() | ||
| _patches_applied = True |
| alias_file = _CONFIG_DIR / f"{device_name}.yaml" | ||
| if alias_file.exists(): | ||
| return alias_file | ||
| except Exception: |
Comment on lines
+7683
to
+7685
| from vllm_fl.dispatch.backends.vendor.gcu.sampler import ( # noqa: F401 | ||
| _random_sample_gcu, | ||
| ) |
| from vllm_fl.dispatch.backends.vendor.gcu.sampler import ( # noqa: F401 | ||
| _random_sample_gcu, | ||
| ) | ||
| except Exception: |
Port the GCU300 enablement to the 0.24.0 dispatch framework (squash of the five-branch history 3402afc..369a551 into one unit): - register enflame in VENDOR_DEVICE_MAP so the platform plugin activates on torch.gcu - add the gcu arm to compilation-graph selection - GCU300 enablement bundle: vendor:gcu attention backend, int64 factory/index blacklist (zeros/add/sub -> torch_gcu), native FLASH_ATTN path wiring - drop generator-seeded exponential from the GCU sampler (correctness) - clamp max_num_batched_tokens to 2047 on GCU Deliberately NOT included here: the attention/operator layer that 0.20.2 delivered (flash_attn_backend binding, register_ops, activation & friends) — that follows as the next commit on top of this one.
tengqm
force-pushed
the
port/enflame-gcu300-v024
branch
from
September 3, 2026 04:53
369a551 to
542026a
Compare
The squash commit carried the platform skeleton (device-map registration, gcu config, slot_mapping int32 rewrite, sampler seed-drop, config clamps) but NOT the operator layer the 0.20.2 enflame run verified — the files patch.py and gcu.py import by name were simply absent, so any dispatch hit (silu_and_mul / rms_norm / rotary_embedding) raised ImportError and the native-FLASH_ATTN binding never ran. Restore from the verified 0.20.2 tree (3431992, vllm-plugin-FL flagos-ai#357); the vLLM internals they patch (v1 attention fa_utils/flash_attn, model_executor.layers.fla.ops, triton_utils) are unchanged in 0.24.0 (metax ships the same fla.ops imports): - impl/activation.py, normalization.py, rotary.py — pure-torch reference impls for the dispatch points gcu.py declares (the missing-import bug) - impl/flash_attn_backend.py — the native-FLASH_ATTN enablement: binds vendor flash_attn.vllm_flash_attn + flag_gems reshape_and_cache_flash onto fa_utils and forces is_flash_attn_varlen_func_available() True (empty wheel strips vllm._C). This is the garbage-decode root cause. - impl/bilinear_pos_embed.py, chunk_delta_h.py, fused_recurrent_packed_decode.py — GCU grid-cap kernel patches (qwen3-vl / FLA linear-attention paths) - patch.py: re-add apply_flash_attn_backend_gcu_patch to apply_gcu_patches - platform.py import_kernels: add the device_type == 'gcu' branch so apply_gcu_patches() actually runs (nothing called it before — same shape as the existing musa branch)
| p_h0_3 = tl.make_block_ptr( | ||
| h0, (V, K), (K, 1), (i_v * BV, 128), (BV, 64), (1, 0) | ||
| ) | ||
| b_h3 += tl.load(p_h0_3, boundary_check=(0, 1)).to(tl.float32) |
| p_h0_4 = tl.make_block_ptr( | ||
| h0, (V, K), (K, 1), (i_v * BV, 192), (BV, 64), (1, 0) | ||
| ) | ||
| b_h4 += tl.load(p_h0_4, boundary_check=(0, 1)).to(tl.float32) |
| chunk_indices: torch.Tensor | None = None, | ||
| chunk_offsets: torch.Tensor | None = None, | ||
| ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor | None]: | ||
| B, T, Hg, K, V = *k.shape, u.shape[-1] |
|
|
||
| qwen3_vl.triton_pos_embed_interpolate = triton_pos_embed_interpolate_gcu | ||
| qwen3_vl._bilinear_pos_embed_kernel = _bilinear_pos_embed_kernel_gcu | ||
| _patched = True |
| # chunk.py / kda.py bind chunk_gated_delta_rule_fwd_h at import time. | ||
| chunk_mod.chunk_gated_delta_rule_fwd_h = chunk_gated_delta_rule_fwd_h_gcu | ||
| kda_mod.chunk_gated_delta_rule_fwd_h = chunk_gated_delta_rule_fwd_h_gcu | ||
| _patched = True |
| gdn_linear_attn.fused_recurrent_gated_delta_rule_packed_decode = ( | ||
| fused_recurrent_gated_delta_rule_packed_decode_gcu | ||
| ) | ||
| _patched = True |
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
2 participants
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.
Summary
Port the Enflame GCU300 enablement bundle to
mainfor vLLM 0.24.0. The twin PRon the
release/0.2line is #357.Changes (1 commit, base
main@ba52028):gcu.py—GCUBackenddispatchers for silu_and_mul / rms_norm /rotary_embedding, plus an attention selector that falls back to the plugin
AttentionFLBackendwhen the empty vLLM wheel ships no compiledflash_attn.vllm_flash_attn._vllm_fa2_C(and aliasessys.modules["vllm.vllm_flash_attn"]when the native path is taken).gcu/impl/slot_mapping.py— on-device int32 slot mapping built withsearchsortedinstead ofrepeat_interleave(grid.y 255 cap).gcu/patch.py— patch module.dispatch/config/gcu.yaml— flag_gems blacklist for the GCU300 int64codegen wall (sort / sort_stable / rsub_scalar / rsub_tensor) and argmax
correctness.
dispatch/config/utils.py— device-name fallback inget_config_pathsogcu.yamlloads forvendor_name == "enflame".Notes
forward), which is closed unmerged on this repo; on
maintheforward_includes_kv_cache_updateattribute defaults toTrueon vLLM'sAttentionBackend, so the read site inmodel_runneris safe. Whether theGCU300 KV-cache path needs the split is to be confirmed on-node.
the 0.20.2 line) and is blacklisted in
gcu.yaml; the Triton path is theverified route.
Verification
Pending on-node: tops1.10.6 and tops1.9.10, app image serve E2E with Qwen3-4B.
This PR was written in part with the assistance of generative AI.