Repository navigation
Conversation
The HIP backend was locked to gfx1100 (RX 7900 XT/XTX) in CMake, the device check, the dp4a intrinsic, and setup.py. gfx1101 and gfx1102 are the same RDNA3 wave32 ISA with the same 64 KiB LDS and the sudot4 dot instruction, so widen those checks to gfx1100/gfx1101/gfx1102. Also guard the native_mmvq_multi test target: its source file is not in the repository, so configuring with STRATA_BUILD_TESTS=ON failed. Tested on an RX 7600 (gfx1102, 8 GB) with system ROCm; gfx1101 is untested. Co-Authored-By: Claude Opus 5.5 <noreply@anthropic.com>
Owner
|
Nice and minimal. One request: compare the runtime arch with the one the binary was compiled for, not the whole RDNA3 family, or a gfx1100 build carried to a gfx1102 card fails later with 'invalid device function'. Please post the end-to-end RX 7600 run when the model is downloaded. |
Owner
|
Thanks @nexus2905. We can't test AMD right now, so closing for now. Please open a new PR when (1) the startup check compares the running card's arch with the one the binary was built for (not the whole RDNA3 family), and (2) you've posted an end-to-end run on the RX 7600. 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. |
shiftySenpai
pushed a commit
to shiftySenpai/Strata
that referenced
this pull request
Sep 30, 2026
…d with the compiled archs; setup picks ROCm and the hipBLASLt table per arch - CMake: CMAKE_HIP_ARCHITECTURES may list gfx1100 and gfx1201 (validated) and gfx1101 / gfx1102 / gfx1200 (build with a warning: same wave32, 64 KiB LDS and dot4, not run by us; Niko1221#192, Niko1221#176). The list reaches the code as STRATA_HIP_ARCHS. - Runtime: the engine (every GPU a run uses) and strata-device compare the card's gcnArchName (up to the ':' feature suffix) with the compiled archs and require wave32, with a message naming the card, its arch and the build's archs, instead of a later "invalid device function". - __dp4a uses __builtin_amdgcn_sudot4 (v_dot4_i32_iu8) on all five archs; the portable fallback stays for other compilers/targets. - setup.py: AMD_ARCHS gfx1100 + gfx1201; TheRock's wheel index per GPU family (gfx110X-dgpu, gfx120X-all; the pinned 7.10.0a20251120 is on both), only that family's libraries on the path, reinstall when the family changes; a system ROCm older than 7.0 is passed over for the wheels; the engine is rebuilt when the chosen card's arch is not in its BUILD.json. The hipBLASLt table is tools/hip/<arch>-hipblaslt-<version>.txt for the card's arch AND the installed hipBLASLt (read from hipblaslt-version.h, 1.4.1 -> 100401), else none with a message; before, the highest-named table of the arch was taken whatever the installed version, which the engine then refused. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Niko1221
added a commit
that referenced
this pull request
Oct 1, 2026
…256 by Efeisot) as community-validated cards Cut down from #254 and #256 onto 0.1.30's AMD tables: - setup.py: AMD_ARCHS / AMD_NAMES / ROCM_INDEXES / AMD_CARDS list gfx1101 (TheRock gfx110X-dgpu) and gfx1200 (gfx120X-all); amd_gpus() takes the sysfs root so the detection can be tested. - cmake/hip_backend.cmake: gfx1101 and gfx1200 build without the "not validated" warning (a status line names them community-validated); gfx1102 (#192: ctest only) keeps the warning. - docs/AMD_HIP.md (Community-validated cards: the owners' setups, ctest and numbers) and README. - tools/test_setup_amd.py: detection on a mocked KFD topology (every family, the CPU node, an iGPU and gfx1102 listed as not supported, product-name fallbacks, the ROCm index per family). Left out: the PRs' own cmake/device.cu arch checks (0.1.30's compiled_for covers them), the dp4a change (main has it), math_constants.h, the bench JSONs, the launcher, and both hipBLASLt tables (neither PR measured the prompt speed without its table). Co-Authored-By: Joseph Hohertz <jhohertz@gmail.com> Co-Authored-By: efeisot <efeproofkmr@gmail.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
Niko1221
added a commit
that referenced
this pull request
Oct 6, 2026
…256 by Efeisot) as community-validated cards Cut down from #254 and #256 onto 0.1.30's AMD tables: - setup.py: AMD_ARCHS / AMD_NAMES / ROCM_INDEXES / AMD_CARDS list gfx1101 (TheRock gfx110X-dgpu) and gfx1200 (gfx120X-all); amd_gpus() takes the sysfs root so the detection can be tested. - cmake/hip_backend.cmake: gfx1101 and gfx1200 build without the "not validated" warning (a status line names them community-validated); gfx1102 (#192: ctest only) keeps the warning. - docs/AMD_HIP.md (Community-validated cards: the owners' setups, ctest and numbers) and README. - tools/test_setup_amd.py: detection on a mocked KFD topology (every family, the CPU node, an iGPU and gfx1102 listed as not supported, product-name fallbacks, the ROCm index per family). Left out: the PRs' own cmake/device.cu arch checks (0.1.30's compiled_for covers them), the dp4a change (main has it), math_constants.h, the bench JSONs, the launcher, and both hipBLASLt tables (neither PR measured the prompt speed without its table). Co-Authored-By: Joseph Hohertz <jhohertz@gmail.com> Co-Authored-By: efeisot <efeproofkmr@gmail.com>
Hugua700
pushed a commit
to Hugua700/Strata
that referenced
this pull request
Oct 6, 2026
…d with the compiled archs; setup picks ROCm and the hipBLASLt table per arch - CMake: CMAKE_HIP_ARCHITECTURES may list gfx1100 and gfx1201 (validated) and gfx1101 / gfx1102 / gfx1200 (build with a warning: same wave32, 64 KiB LDS and dot4, not run by us; Niko1221#192, Niko1221#176). The list reaches the code as STRATA_HIP_ARCHS. - Runtime: the engine (every GPU a run uses) and strata-device compare the card's gcnArchName (up to the ':' feature suffix) with the compiled archs and require wave32, with a message naming the card, its arch and the build's archs, instead of a later "invalid device function". - __dp4a uses __builtin_amdgcn_sudot4 (v_dot4_i32_iu8) on all five archs; the portable fallback stays for other compilers/targets. - setup.py: AMD_ARCHS gfx1100 + gfx1201; TheRock's wheel index per GPU family (gfx110X-dgpu, gfx120X-all; the pinned 7.10.0a20251120 is on both), only that family's libraries on the path, reinstall when the family changes; a system ROCm older than 7.0 is passed over for the wheels; the engine is rebuilt when the chosen card's arch is not in its BUILD.json. The hipBLASLt table is tools/hip/<arch>-hipblaslt-<version>.txt for the card's arch AND the installed hipBLASLt (read from hipblaslt-version.h, 1.4.1 -> 100401), else none with a message; before, the highest-named table of the arch was taken whatever the installed version, which the engine then refused.
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…d with the compiled archs; setup picks ROCm and the hipBLASLt table per arch - CMake: CMAKE_HIP_ARCHITECTURES may list gfx1100 and gfx1201 (validated) and gfx1101 / gfx1102 / gfx1200 (build with a warning: same wave32, 64 KiB LDS and dot4, not run by us; Niko1221#192, Niko1221#176). The list reaches the code as STRATA_HIP_ARCHS. - Runtime: the engine (every GPU a run uses) and strata-device compare the card's gcnArchName (up to the ':' feature suffix) with the compiled archs and require wave32, with a message naming the card, its arch and the build's archs, instead of a later "invalid device function". - __dp4a uses __builtin_amdgcn_sudot4 (v_dot4_i32_iu8) on all five archs; the portable fallback stays for other compilers/targets. - setup.py: AMD_ARCHS gfx1100 + gfx1201; TheRock's wheel index per GPU family (gfx110X-dgpu, gfx120X-all; the pinned 7.10.0a20251120 is on both), only that family's libraries on the path, reinstall when the family changes; a system ROCm older than 7.0 is passed over for the wheels; the engine is rebuilt when the chosen card's arch is not in its BUILD.json. The hipBLASLt table is tools/hip/<arch>-hipblaslt-<version>.txt for the card's arch AND the installed hipBLASLt (read from hipblaslt-version.h, 1.4.1 -> 100401), else none with a message; before, the highest-named table of the arch was taken whatever the installed version, which the engine then refused.
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…060 XT, Niko1221#256 by Efeisot) as community-validated cards Cut down from Niko1221#254 and Niko1221#256 onto 0.1.30's AMD tables: - setup.py: AMD_ARCHS / AMD_NAMES / ROCM_INDEXES / AMD_CARDS list gfx1101 (TheRock gfx110X-dgpu) and gfx1200 (gfx120X-all); amd_gpus() takes the sysfs root so the detection can be tested. - cmake/hip_backend.cmake: gfx1101 and gfx1200 build without the "not validated" warning (a status line names them community-validated); gfx1102 (Niko1221#192: ctest only) keeps the warning. - docs/AMD_HIP.md (Community-validated cards: the owners' setups, ctest and numbers) and README. - tools/test_setup_amd.py: detection on a mocked KFD topology (every family, the CPU node, an iGPU and gfx1102 listed as not supported, product-name fallbacks, the ROCm index per family). Left out: the PRs' own cmake/device.cu arch checks (0.1.30's compiled_for covers them), the dp4a change (main has it), math_constants.h, the bench JSONs, the launcher, and both hipBLASLt tables (neither PR measured the prompt speed without its table). Co-Authored-By: Joseph Hohertz <jhohertz@gmail.com> Co-Authored-By: efeisot <efeproofkmr@gmail.com>
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…d with the compiled archs; setup picks ROCm and the hipBLASLt table per arch - CMake: CMAKE_HIP_ARCHITECTURES may list gfx1100 and gfx1201 (validated) and gfx1101 / gfx1102 / gfx1200 (build with a warning: same wave32, 64 KiB LDS and dot4, not run by us; Niko1221#192, Niko1221#176). The list reaches the code as STRATA_HIP_ARCHS. - Runtime: the engine (every GPU a run uses) and strata-device compare the card's gcnArchName (up to the ':' feature suffix) with the compiled archs and require wave32, with a message naming the card, its arch and the build's archs, instead of a later "invalid device function". - __dp4a uses __builtin_amdgcn_sudot4 (v_dot4_i32_iu8) on all five archs; the portable fallback stays for other compilers/targets. - setup.py: AMD_ARCHS gfx1100 + gfx1201; TheRock's wheel index per GPU family (gfx110X-dgpu, gfx120X-all; the pinned 7.10.0a20251120 is on both), only that family's libraries on the path, reinstall when the family changes; a system ROCm older than 7.0 is passed over for the wheels; the engine is rebuilt when the chosen card's arch is not in its BUILD.json. The hipBLASLt table is tools/hip/<arch>-hipblaslt-<version>.txt for the card's arch AND the installed hipBLASLt (read from hipblaslt-version.h, 1.4.1 -> 100401), else none with a message; before, the highest-named table of the arch was taken whatever the installed version, which the engine then refused. Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
anon761
pushed a commit
to anon761/Strata-upstream
that referenced
this pull request
Oct 6, 2026
…060 XT, Niko1221#256 by Efeisot) as community-validated cards Cut down from Niko1221#254 and Niko1221#256 onto 0.1.30's AMD tables: - setup.py: AMD_ARCHS / AMD_NAMES / ROCM_INDEXES / AMD_CARDS list gfx1101 (TheRock gfx110X-dgpu) and gfx1200 (gfx120X-all); amd_gpus() takes the sysfs root so the detection can be tested. - cmake/hip_backend.cmake: gfx1101 and gfx1200 build without the "not validated" warning (a status line names them community-validated); gfx1102 (Niko1221#192: ctest only) keeps the warning. - docs/AMD_HIP.md (Community-validated cards: the owners' setups, ctest and numbers) and README. - tools/test_setup_amd.py: detection on a mocked KFD topology (every family, the CPU node, an iGPU and gfx1102 listed as not supported, product-name fallbacks, the ROCm index per family). Left out: the PRs' own cmake/device.cu arch checks (0.1.30's compiled_for covers them), the dp4a change (main has it), math_constants.h, the bench JSONs, the launcher, and both hipBLASLt tables (neither PR measured the prompt speed without its table). Co-Authored-By: Joseph Hohertz <jhohertz@gmail.com> Co-Authored-By: efeisot <efeproofkmr@gmail.com> Co-Authored-By: Claude Opus 5.5 (1M context) <noreply@anthropic.com> Claude-Session: https://claude.ai/code/session_0154bwiad8gY7wRehCXy27pN
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
2 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.
What
The experimental HIP backend was locked to gfx1100 (RX 7900 XT/XTX) in four places. This PR widens those checks to also accept gfx1101 and gfx1102, the rest of the RDNA3 wave32 family:
cmake/hip_backend.cmake: acceptCMAKE_HIP_ARCHITECTURESmatchinggfx110[012]src/core/device.cu: the runtime device check accepts gfx1100/1101/1102 (still requires wave32)include/strata/hip_compat/intrinsics.hpp: use__builtin_amdgcn_sudot4fordp4aon gfx1101/gfx1102 too, not just the scalar fallbacksetup.py:AMD_ARCHS/AMD_NAMESlist the RX 7700/7800 (gfx1101) and RX 7600 (gfx1102), and the "not supported" messages are updated to matchsrc/core/device_main.cpp:strata-deviceprints "RDNA3 wave32" instead of a hard-coded gfx1100Unrelated fix:
CMakeLists.txtguards thenative_mmvq_multitest target withEXISTS.bench/micro/native_mmvq_multi.cppis missing from the repository, so configuring withSTRATA_BUILD_TESTS=ONfailed.Why
gfx1101 and gfx1102 use the same RDNA3 wave32 ISA as gfx1100. They have the same 64 KiB LDS limit that
fused_gr.cu's tile size assumes, and they also have the signed dot4 instruction. Nothing in the kernels depends on gfx1100 specifically. The one hardware difference is that gfx1102 has a smaller VGPR file, which affects occupancy, not correctness. The hipBLASLt tuning table intools/hipis keyed by arch, so on these cards setup doesn't find a table and the engine falls back to hipBLASEx, as it already does when a table doesn't match.Testing
Tested only on an RX 7600 (gfx1102, 8 GB) + Ryzen 5 5600 (AVX2, no AVX-512) + 64 GB RAM, Ubuntu 26.04, system ROCm in
/opt/rocm:stratabuilds with-DCMAKE_HIP_ARCHITECTURES=gfx1102, both by hand and through./setup.sh --backend hip, which detects the card and compiles the enginestrata-devicereports the card correctly and plans an automatically sized expert cache (about 4,100 slots in 8 GB)ctest(excludingple_parityandplatform_memory_test, as in docs/AMD_HIP.md): 39/40 pass, including every HIP parity test (intrinsics, handoff, native QSA, GR, quantize_act, KV, MMQ prefill, ...).hip_prefill_hipblaslt_gemmwas skippedexpert_multi_test: it exercises the AVX-512 CPU expert kernel and exits on a CPU without AVX-512. That isn't related to this change; the engine chooses the AVX2 kernels on such CPUsNot tested: gfx1101 (RX 7700/7800) on real hardware, and a full end-to-end model run on gfx1102. The model download for that is still in progress. I haven't touched docs/AMD_HIP.md; I can update it if you want these cards listed there.
🤖 Generated with Claude Code