Skip to content

hip: gfx1100 hipBLASLt tuning table for ROCm 7.2.1 (100202) - #170

Closed
samuelishida wants to merge 1 commit into
Niko1221:mainfrom
samuelishida:pr-hipblaslt-100202
Closed

samuelishida wants to merge 1 commit into
Niko1221:mainfrom
samuelishida:pr-hipblaslt-100202

Conversation

@samuelishida

Copy link
Copy Markdown
Contributor

Adds tools/hip/gfx1100-hipblaslt-100202.txt, calibrated with the branch's own tuner on an RX 7900 XTX under ROCm 7.2.1.

  • 26 of the 30 covered shapes validated.
  • Four shapes at T=4096 produce non-finite output in every hipBLASLt solution on 100202; those rows are intentionally absent and the runtime falls back to hipblasGemmEx.

Also makes hip_prefill_hipblaslt_gemm self-locate a matching shipped table for the runtime's hipBLASLt version (falling back from STRATA_HIPBLASLT_TUNING), so the parity test runs unconditionally on a machine with a committed table. Documents the table and its fallback behavior in docs/AMD_HIP_PERFORMANCE.md.

Adds tools/hip/gfx1100-hipblaslt-100202.txt, calibrated with the branch's own
tuner on an RX 7900 XTX. 26 of the 30 covered shapes validated; four shapes at
T=4096 produce non-finite output in every hipBLASLt solution on 100202, so those
rows are intentionally absent and the runtime falls back to hipblasGemmEx.

Makes hip_prefill_hipblaslt_gemm self-locate a matching shipped table for the
runtime's hipBLASLt version, falling back from STRATA_HIPBLASLT_TUNING so the
parity test runs unconditionally on a machine with a committed table.
Documents the 100202 table and its fallback behavior in docs/AMD_HIP_PERFORMANCE.md.
@Niko1221

Copy link
Copy Markdown
Owner

Thanks @samuelishida. Two things for a new PR: the test uses STRATA_TUNING_TABLE_DIR, which isn't defined in CMake (the HIP test build breaks), and the doc implies the engine finds the table itself; only the test does. Also note our setup currently takes the highest-numbered gfx1100 table regardless of the installed hipBLASLt version; we'll fix setup to pick the table that matches the runtime, and then a 100202 table can go in cleanly. We're closing PRs that wait on changes so the list stays manageable; this is not a no. Please open a new PR when it's ready and link this one.

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

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants