feat(ds4): opt-in GPU F16 HC fn mirrors for ROCmFPX#546
Draft
Graffioh wants to merge 4 commits into
Draft
Conversation
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce
Graffioh
force-pushed
the
codex/ds4-rocmfpx-hc-gpu-mirror
branch
from
July 19, 2026 09:28
1c69b8a to
8d59290
Compare
Also size the shared scratch for the requested shapes: the bounded buffers (mix/post/comb) are allocated at their kMaxHc maxima and the working buffer tracks n_embd capacity, instead of staying at whatever dims the first call happened to use. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce
ROCmFPX-quantized HC fn tensors cannot feed the __half HC pre kernels, so HC pre always fell back to CPU on ROCmFPX builds. Behind DFLASH_DS4_ROCMFPX_HC_GPU (default off), upload the existing host F16 copy of each in-range layer's attn/ffn fn weights once per shard cache and let the GPU HC pre path consume it. Mirrors are owned by the shard's DeepSeek4LayerRangeCache, keyed to the owning device, and freed on that device. A failed upload logs an explicit CPU-HC fallback and degrades gracefully. The output HC projection is not mirrored: it runs once per step, not per layer. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce
Graffioh
force-pushed
the
codex/ds4-rocmfpx-hc-gpu-mirror
branch
from
July 19, 2026 09:53
8d59290 to
b47e409
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
1 participant
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
Opt-in performance feature for ROCmFPX builds: GPU F16 mirrors for quantized HC fn weights, gated behind
DFLASH_DS4_ROCMFPX_HC_GPU(default off).ROCmFPX-quantized HC
fntensors cannot feed the__halfHC pre kernels, so HC pre falls back to CPU on every layer, every token. With the flag on, the existing host F16 copy of each in-range layer's attn/ffn fn weights is uploaded once per shard cache and the GPU HC pre path consumes it.Stack
Implementation
deepseek4_cuda_hc_upload_f16/deepseek4_cuda_hc_freedevice-memory primitives (select device, malloc + H2D upload; best-effort free on the owning device)HcWeightsCpu(fn_f16_device+ size + owning device id), uploaded ininitialize_layer_range_cachefor the shard's layer range onlyhc_fn_device_ptrreturnsfn->datafor F16 tensors (unchanged), the mirror for quantized tensors when the flag is on, else nullptr → CPU HC per layerValidation / TODO before ready
hc_dim x mix_dimF16, x2 for attn + ffn) still required