455 war room findings #1001
Open
jpvillam-amd wants to merge 23 commits into
Open
Conversation
Squashed application of rocm/vllm gfx1250_wip_dllehr (19 commits, 3 internal merges) onto current origin/main (~2666 commits newer). Conflict resolutions: - docker/Dockerfile.rocm, docker/Dockerfile.rocm_base, .buildkite/ release-pipeline.yaml, .buildkite/scripts/annotate-rocm-release.sh: kept WIP versions (TheRock-based ROCm build automation). - CMakeLists.txt: upstream HIP_SUPPORTED_ARCHS list + gfx1250. - vllm/platforms/rocm.py: kept upstream _ON_GFX12X/_ON_GFX90A AND added gfx1250 to _ON_MI3XX/_ON_GFX9 (WIP intent). - vllm/model_executor/layers/quantization/utils/mxfp4_utils.py: re-applied the live gfx1250 scale_layout change (on_gfx950() or on_gfx1250()); dropped the stale _can_support_mxfp4/get_padding_alignment (upstream removed them + callers). - vllm/_aiter_ops.py (new), vllm/envs.py, vllm/v1/worker/gpu_worker.py, csrc/quickreduce/base.h: applied cleanly. - vllm/model_executor/layers/attention/mla_attention.py: resolved to upstream; the WIP's ROCm flash_attn fallback was relocated upstream to vllm/v1/attention/backends/fa_utils.py. DEFERRED: porting the aiter.ops.triton.mha.flash_attn_varlen_func substitution there. - vllm/model_executor/layers/quantization/quark/quark_moe.py: resolved to upstream (it was rewritten +747/-304 into a backend/kernel-factory model). DEFERRED: the VLLM_ROCM_AITER_FUSED_MOE_TRITON_GEMM_A4W4 shuffle-skip guardrail now belongs in fused_moe/oracle/mxfp4.py (AITER_MXFP4_MXFP4 W4A4 backend selection), an area being reworked separately. AI-assisted (Claude Code); human review required before any PR, especially the two DEFERRED ports above. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Completes deferred port #1 from the gfx1250 squash. FlashAttention (the 'flash_attn' pip package) is not installed on ROCm, so fa_utils.py's ROCm branch now imports flash_attn_varlen_func from aiter.ops.triton.mha (same source as aiter_triton_mla.py) instead of 'flash_attn'. _ROCM_FLASH_ATTN_AVAILABLE is set True when AITER is present, so is_flash_attn_varlen_func_available() reports a working impl. All ROCm consumers that import flash_attn_varlen_func from fa_utils (mla/prefill/flash_attn.py, flash_attn_diffkv.py, turboquant_attn.py, the vit_attn_wrappers fallback) inherit this. The only other 'import flash_attn' (qwen2_5_omni_thinker.py) is already guarded (-> None). Verified in the gfx1250 FFM image: resolves to aiter.ops.triton.attention.mha.flash_attn_varlen_func, availability True. AI-assisted (Claude Code); human review + lint pending before any PR. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Enables a from-source vLLM build with PYTORCH_ROCM_ARCH=gfx1250 (verified in the
FFM gfx1250 simulator image; _C and _rocm_C carry gfx1250 code objects, import OK):
- csrc/rocm/attention.cu: guard the gfx12 WMMA builtins
(__builtin_amdgcn_wmma_f32_16x16x16_{f16,bf16}_w32_gfx12) for gfx1250, which
lacks them (needs wmma-128b-insts); trap if launched.
- CMakeLists.txt + csrc/rocm/torch_bindings.cpp: skinny_gemms.cu is pervasively
gfx9/gfx11 ISA (MFMA, dot2/dot4, legacy s_waitcnt asm) unsupported on gfx1250.
When VLLM_GPU_ARCHES matches gfx1250, exclude skinny_gemms.cu from _rocm_C and
define VLLM_SKIP_SKINNY_GEMMS to skip its op registrations (LLMM1/wvSplitK/
wvSplitKrc/wvSplitKQ). vLLM falls back to default/Triton GEMM for these on
gfx1250. Non-gfx1250 ROCm builds are unaffected.
Caveat: these custom ROCm GEMM ops are unavailable on gfx1250 and the custom
attention WMMA path traps; intended for functional bring-up on the simulator.
AI-assisted (Claude Code); human review + real gfx1250 run needed before any PR.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Lets vLLM initialize and run on the gfx1250 FFM pre-silicon simulator (verified: facebook/opt-125m generates end-to-end with attention_backend=TRITON_ATTN): - platforms/__init__.py: ROCm platform detection falls back to torch.version.hip when amdsmi is unavailable. amdsmi queries the driver/sysfs and does not see GPUs exposed only via HSA (the FFM model); torch's HIP runtime does. - platforms/rocm.py: _get_gcn_arch() uses logger.debug (not warning_once) on the amdsmi-failure path; warning_once imports vllm.distributed at module load, causing a circular import before current_platform is bound. The torch.cuda gcnArchName fallback then correctly reports gfx1250 (amdsmi would report the host's real gfx950 cards, not the sim). - model_executor/layers/utils.py: skip the skinny-GEMM (wvSplitK/LLMM1) path on gfx1250 (those kernels are excluded from _rocm_C); fall back to torch GEMM. Note: on gfx1250 the ROCm custom-attention backends route paged decode through _rocm_C.paged_attention; use attention_backend=TRITON_ATTN (pure Triton) for now. AI-assisted (Claude Code); human review needed before any PR. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…er-prune loader More fixes to bring vLLM up on the gfx1250 FFM simulator (no amdsmi; amdsmi would report the host's real gfx950 cards, not the sim): - platforms/rocm.py: add _AMDSMI_AVAILABLE flag; with_amdsmi_context skips amdsmi init/shutdown when unavailable; get_device_name falls back to torch.cuda.get_device_name. (Complements the earlier _get_gcn_arch fix.) - models/deepseek_v4/amd/model.py: load_weights skips checkpoint weights for transformer layers pruned via a num_hidden_layers hf_override (was raising KeyError), enabling small-layer-count bring-up runs. Probe result (DeepSeek-V4-Flash, 4 layers, TP=1, AITER off, triton_unfused MoE): loads and runs the FP8 attention via the Triton GEMM, but the MXFP4 MoE fails to compile -- triton_kernels matmul_ogs uses unswizzle_mx_scale_cdna4 (gfx950 scale layout) for gfx1250; reshape mismatch. gfx1250 needs its own MXFP4 scale-unswizzle (cf. aiter swizzle_scales_gfx1250). MoE is the gfx1250 enablement gap. AI-assisted (Claude Code); human review needed. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
… on gfx1250 Routes the DeepSeek-V4-Flash MXFP4 MoE through aiter's gfx1250 W4A8 `moe_gemm_a8w4` instead of the vendored triton_kernels `matmul_ogs`, whose `unswizzle_mx_scale_cdna4` has no gfx1250 variant. - gpt_oss_triton_kernels_moe.py: add `UnfusedOAITritonExperts._try_apply_aiter_w4a8` plus an early-return gate in `apply()` (fires on use_mxfp4_w4a8/w4a16, no router-weight-on-input, no LoRA, rocm_aiter_ops enabled). Does a manual expert-sorted gather (in-kernel gather is numerically broken on gfx1250), builds aiter-native routing from topk_ids/topk_weights, resolves dynamic FP8 LHS scales, runs gemm1(swiglu, fp8 out) -> gemm2, and restores DeepSeek-V4's routed_scaling_factor via a per-token output rescale. - mxfp4_utils.py: store a plain StridedLayout MXFP4 scale on gfx1250 (only gfx950 keeps the CDNA4/GFX950 swizzle); the W4A8 path reads it un-swizzled. - deepseek_v4/compressor.py: on ROCm, select the pure-Triton sparse-attention compressor (head_dim==512) instead of the NVIDIA-only CuTeDSL path. - test_modular_oai_triton_moe.py: coverage for the W4A8 MoE path. Verified end-to-end on the gfx1250 FFM simulator: DeepSeek-V4-Flash loads and runs MLA + sparse attention + MXFP4 MoE to generation. MoE kernel math validated standalone to maxrel ~5e-3. Diagnostic logger.info calls in `_try_apply_aiter_w4a8` are retained intentionally for sim bring-up; drop them before an upstream PR. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Signed-off-by: Douglas Lehr <douglehr@amd.com>
…oaders
Enables `hf_overrides={"num_hidden_layers": N}` layer-prune smoke tests for
gpt-oss-120b on the FFM sim (mirrors the DSV4 loader fix). Both
`_load_weights_mxfp4` and `_load_weights_quark` did raw `params_dict[name]`
lookups and KeyError'd on checkpoint weights for layers the pruned model
doesn't build. Skip any weight whose `layers.<idx>.` index is >=
num_hidden_layers.
WIP bring-up: with this, gpt-oss-120b (mxfp4) loads but the forward still
crashes on gfx1250 with an async "invalid device function" upstream of the MoE
experts (kernel not pinpointed yet). The AMD quark w4a8 variant
(/data/amd/gpt-oss-120b-w-mxfp4-a-fp8) is staged to try next.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Signed-off-by: Douglas Lehr <douglehr@amd.com>
… path) Generalizes the monolithic AITER MXFP4 W4A8 MoE experts class so it can fire on gfx1250 (RDNA4 sim), not just gfx950. This is the code path gpt-oss-120b selects when use_mxfp4_w4a8 is set -- the earlier DSV4 commit only enabled the gpt_oss_triton_kernels path, not AiterW4A8ExpertsMonolithic. aiter_triton_kernel_w4a8_moe_forward / triton_kernel_fused_mxfp4_w4a8_experts: - Routing: on gfx1250 import aiter's pure-torch `routing_torch` instead of the triton `routing`. The triton routing kernel compiles a TMA (TDM) descriptor whose last dim is `topk * 2` bytes, which is < 16 (the descriptor minimum) for power-of-2 topk like gpt-oss' topk=4, and fails to compile for the warmup dummy run plus every decode step. routing_torch is numerically identical to routing on the sim where the latter compiles. DSV4's topk=6 dodged this because next_power_of_2(6)==8!=6 disables the descriptor branch in aiter. - Try aiter's nested `moe.moe_routing.routing` module first, fall back to the legacy `moe_routing.routing` path (handles aiter version skew). - Gather: gfx1250 in-kernel gather is numerically broken (validated on the sim: do_gather=True -> maxrel ~2.4). Replicate aiter's gather in torch (sorted row i reads token gather_idx[i] // topk) and pass gather_indx=None; reproduces the in-kernel gather to ~5e-3 maxrel. gfx950 path unchanged. - Scale swizzle: gfx1250 stores the MXFP4 weight scale unswizzled (mxfp4_utils._swizzle_mxfp4 keeps StridedLayout on gfx1250), and the gfx1250 moe_gemm_a8w4 reads CDNA4_SCALE as garbage (CDNA4_SCALE -> maxrel ~7e4 on the sim, None -> ~6e-3). Pass swizzle_mx_scale=None on gfx1250; gfx950 still uses "CDNA4_SCALE". _supports_current_device(): accept on_gfx950() OR on_gfx1250(). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Signed-off-by: Douglas Lehr <douglehr@amd.com>
- add libpciaccess0/libpciaccess-dev/libdrm-dev/ pkg-config/cmake to base apt; drop the standalone cmake install - install rocm[libraries,devel] (vs rocm[devel]) and pin torch/torchvision/torchaudio via genesis gfx1250 index - /opt/rocm symlinks for bin/include/lib and rocprofiler-sdk; PATH and LD_LIBRARY_PATH wired through. - AITER: switch branch shared/triton-gfx12 -> main. - ENABLE_CK toggle: AITER_ENABLE_CK=0 (default) disables Composable Kernel for the current build Signed-off-by: Daniel <danichan@amd.com>
* Add 1250 only base+vllm Dockerfile * Add upstream dockerfiles for rocm and rocm base * Enable rocm_base for gfx1250 --------- Co-authored-by: jpvillam <jpvillam@amd.com>
Signed-off-by: jpvillam <juan.villamizar@amd.com>
Signed-off-by: jpvillam <juan.villamizar@amd.com>
Signed-off-by: jpvillam <juan.villamizar@amd.com>
Signed-off-by: jpvillam <juan.villamizar@amd.com>
Signed-off-by: jpvillam <juan.villamizar@amd.com>
added 2 commits
June 24, 2026 12:50
danichan-mkm
force-pushed
the
455_wip
branch
2 times, most recently
from
e7aa711 to
f18b3f8
Compare
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
3 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.
No description provided.