Skip to content

RX 7900 XTX (gfx1100), Ubuntu, setup's ROCm 7.10 wheel: ~0.8 tok/s decode, GPU at 100% / ~320 W #590

Closed as not planned

Description

@ibrokers77882

AMD RX 7900 XTX (gfx1100): decode ~0.76 tok/s, prefill ~1 tok/s; every GPU stage uniformly slow, raw GPU bandwidth healthy

System: Ubuntu [Ubuntu 26.04 LTS], kernel [7.0.0-34-generic], Strata commit [1678de3] (engine 0.1.34). Ryzen 9 7900X, 96 GB RAM,
RX 7900 XTX 24 GB (gfx1100), display on the iGPU (XTX has ~0 desktop VRAM use). setup.py HIP build with the
TheRock 7.10.0a20251120 gfx110X-dgpu wheels, clang 22, libstdc++-16-dev. PCIe Gen4 x16.

Config: IQ3_S, --expert-cache auto --prefill auto --spec 4 --spec-min-p 0.5 --kv int8 --kv-resident 32768 --max-context 131072 --pcie-frac 0 --mmap-experts --resident-cpu-experts.

Symptom: 58-token prompt takes ~50 s to prefill (~1 tok/s); decode 0.76 tok/s (2.5 tokens/window,
~3300 ms/window). Docs list ~55 tok/s for this card, i.e. ~70x slower. GPU shows 100% busy at 240-390 W
while generating; 0% / 20 W idle.

STRATA_DECODE_TIMING: 3300.82 ms/window = verify 3045 (GPU-reach wait 2695, per-layer host 21 ms) + draft 253.
CPU experts 6.76, VRAM hits 21.46, PCIe 0.00 per layer-window. waitCPU 0.27 ms.

STRATA_VERIFY_PROFILE (ms/window, GDN layers): q8+qkv gemv 489, z 278, out-proj 234, hc-read1+router 323,
shared+quant 185, VRAM hits 516, head 368. Nothing dominates; all stages are slow by a similar factor.

Ruled out:

  • GPU/driver/clocks: standalone HIP copy test on the same card: 687 GB/s (hipMemcpy D2D), 676 GB/s (own kernel,
    built with the same clang 22). mclk at max state, performance level auto, no resets or timeouts in dmesg.
  • Memory mode: same speed in the default pinned-arena mode and with --mmap-experts --resident-cpu-experts.
  • PCIe (stage ~1%, link Gen4 x16), swap and disk I/O, CPU threads, expert hit rate (~85%), VRAM spill
    (same speed with 2.9 GiB free), GTT limit.
  • iGPU: HIP_VISIBLE_DEVICES=0, iGPU idle (0%) during generation.
  • User env: HSA_OVERRIDE_GFX_VERSION=11.0.0 and ROCBLAS_USE_HIPBLASLT=1 were set in my shell; removing both
    (verified absent in the engine's /proc environ) did not change timing (1m49s for 40 tokens either way).
  • Build type: Release. Libraries: engine loads only the wheel's libamdhip64/hsa/comgr/rocblas/hipblaslt.

Could not profile: rocprofv3 --attach fails (couldn't dlsym rocprofiler_register_attach); running the server
under rocprofv3 --kernel-trace segfaults in rocprofiler::hsa::WriteInterceptor during hipGraphLaunch.

Questions: Which ROCm/clang versions did you measure on gfx1100 (docs mention ROCm 7.1.52802 / clang 20)?
Is a newer wheel index for gfx110X known to work? Is there a built-in kernel benchmark I can run to find which kernel is slow?

strata-log-clean.txt

redacted-stacks1.txt

redacted-stacks3.txt

Activity

  1. Niko1221 commented on Oct 3, 2026

    @Niko1221
    Owner

    Thanks for the thorough report. The docs' gfx1100 numbers were measured with ROCm 7.1 (7.1.52802); setup installs the TheRock 7.10.0a20251120 wheel, and your Ubuntu 26.04 / kernel 7.0 combination is new to us. Every stage being slow by the same factor, with healthy copy bandwidth, points at the environment rather than one kernel.

    Could you:

    1. update to 0.1.38 (./update.sh) and retest with the config setup writes (without the extra --pcie-frac 0 --mmap-experts --resident-cpu-experts), then post the first ~80 lines of the log;
    2. check whether another ROCm app (for example llama.cpp's HIP build) reaches its normal speed on the same install?

    If 1 is still slow and 2 is normal, we'll know it's on our side.

  2. ibrokers77882 commented on Oct 3, 2026

    @ibrokers77882
    Author

    Thanks. Both checks done:

    1. Updated to 0.1.38 (commit 99f3dbd, engine rebuilt locally), removed my extra --pcie-frac 0 --mmap-experts --resident-cpu-experts,
      default config from setup. Still slow: 40-token reply took 2m04s (was 1m49s on 0.1.34), 3366 ms/window.
      First 80 log lines attached (log-head80.txt).
    2. llama.cpp (commit 836d571) built with the SAME TheRock 7.10.0a20251120 toolchain (clang 22) and run on the same runtime:
      Qwen2.5-7B Q4_K_M, -ngl 99, RX 7900 XTX: pp512 = 3992 tok/s, tg128 = 134 tok/s. So the ROCm stack/driver/card are fine.

    So it looks like it is on your side. Possible lead: decode ends each layer with the host-mapped doorbell the GPU polls;
    the GPU shows 100% busy at 240-390 W while generating. I tried --gpu-only-full to separate GPU time from the host handshake,
    but a native pack refuses it ("not captured with post graphs"). Is there another way to get a pure-GPU number for a
    native pack, or other flags/env vars you want me to try (e.g. STRATA_POOL_SPIN_US)?

    Extra: RAM runs at 5600 MT/s; Ubuntu 26.04, kernel 7.0.

    log-head80.txt

  3. ibrokers77882 commented on Oct 4, 2026

    @ibrokers77882
    Author

    For comparison on the same machine: llama.cpp (836d571, same wheel toolchain) runs Unsloth UD-IQ3_XXS (82 GB, experts split CPU/GPU with --n-cpu-moe) at ~25 tok/s generation and ~390-418 tok/s prompt processing (llama-bench, 7900 XTX + 96 GB RAM). Strata 0.1.38 on the same box: 0.76 tok/s.

  4. Niko1221 commented on Oct 13, 2026

    @Niko1221
    Owner

    Thanks for the report. To keep the tracker current we are closing issues older than three days: 0.1.44 is out with a long list of fixes, and many older reports are covered by it or by 0.1.43.1.

    If this still happens on 0.1.44, please leave a comment here with your engine log (and the version line it prints), and we will reopen it. A new issue that links this one works too.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Labels

    No labels
    No labels

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions