Skip to content

ROCm 7.2 + RDNA4: performance silently collapses (CK FMHA probe fails, ~2x slower) #1245

Description

@nlzy

Summary

On ROCm 7.2.x with RDNA4 GPU. The build silently falls back to unfused attention.

With transformer nets this costs about half of NN throughput, with no warning at build time or runtime.

Environment

  • KataGo v1.18.1, -DUSE_BACKEND=ROCM -DGPU_TARGETS=gfx1200
  • RX 9060 XT 16 GB (gfx1200), ROCm 7.2.3, Debian GNU/Linux 13
  • Model: kata1-tf3-b11c768-s11001M-d5973M.bin.gz

Performance

Same machine, same config (numSearchThreads=16, 500-visit benchmark):

Build nnEvals/s
Probe failed (ROCm 7.2 headers) 204
CK FMHA (therock-7.13 headers) 434

The CK FMHA build uses -DKATAGO_CK_TILE_INCLUDE_DIR pointing at manually downloaded therock-7.13 headers.

Both paths pass testgpuerror -quick with equal error, so this is purely a performance loss.

Suggested fixes

  1. Warn the user clearly at both build time and runtime.
    • Pros: trivial change; no behavior change, no code risk.
    • Cons: does not restore performance.
  2. Fetch the required headers at CMake time (pinned rocm-libraries tag).
    • Pros: actually fixes the problem; no repo size growth.
    • Cons: build needs network access; depends on an external service.
  3. Vendor the matching ck_tile headers into the KataGo repo.
    • Pros: actually fixes the problem; no network dependency.
    • Cons: repo grows (~26 MB of headers).

Happy to test any patch on gfx1200 / ROCm 7.2.

Activity

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