Add DISABLE_TARGET_PROJECTS for architecture-independent targets - #6959
Draft
idubinov wants to merge 13 commits into
Draft
Add DISABLE_TARGET_PROJECTS for architecture-independent targets#6959idubinov wants to merge 13 commits into
idubinov wants to merge 13 commits into
Conversation
Run black on new_amdgpu_family_matrix.py and add unit tests covering the amdgcnspirv SPIR-V target: build-only/linux-only invariants in the family matrix, and select_targets honoring the non-gfx 'amdgcnspirv' PR label. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The CI matrix (amdgpu_family_matrix.py / new_amdgpu_family_matrix.py)
emits the amdgcnspirv family as "gpu-generic", but the CMake target was
registered as FAMILY generic. The mismatch made THEROCK_DIST_AMDGPU_FAMILIES
validation fail ("value 'gpu-generic' unknown") and broke Python packaging
(expand_families). "generic" is also reserved for architecture-independent
host artifacts, so use gpu-generic consistently.
Also add gpu-generic to the amdgpu_family choice list in
test_linux_jax_wheels.yml so fetch_package_targets_test passes.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Excluded per-arch projects (e.g. all of math-libs/ml-libs for the portable amdgcnspirv target) previously fell back to an unrelated gfx arch, which then failed either the project's own gfx allowlist at configure time (rocBLAS) or the kpack splitter (no device objects matched the requested target). Add therock_project_has_no_targets() and skip those projects outright, generalizing the existing composable_kernel idiom. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The cpp-sdk-user smoke test find_package()s each ROCm math/ML library gated on THEROCK_ENABLE_*. When a library is skipped for the current target (e.g. the portable amdgcnspirv target), those flags are still ON but the package is not installed, so the test fails to configure. Feed the example the actual built status via therock_project_has_no_targets(), matching the existing hipTensor handling. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
libhipcxx is header-only and produces no per-arch device code, so it builds fine for amdgcnspirv and must stay available: hipthreads BUILD_DEPs it. Drop it from the exclude list. hipthreads itself compiles per-arch device bitcode (USE_TEST_AMDGPU_TARGETS), so it belongs in the exclusion and is now skipped via the same gate. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The Python packaging job expands dist families to device targets and, under KPACK_SPLIT_ARTIFACTS, requires a device artifact per target. Architecture- independent targets (the gpu-generic family, e.g. amdgcnspirv) emit no per-arch device code, so they have no device artifact and tripped the completeness check. Filter gpu-generic family members out of the per-device target lists. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
The PyTorch device build compiles CK kernels for every family in its matrix via PYTORCH_ROCM_ARCH. The architecture-independent gpu-generic family (amdgcnspirv) is not a valid device target and fails the compile (gfx-specific intrinsics). Filter it out for all refs/platforms, alongside the per-ref unsupported set. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
EXCLUDE_TARGET_PROJECTS only strips a gfx target from a project's GPU_TARGETS; the project still builds. It cannot express "this project is meaningless for this target and must not build at all", which is exactly what the build-only amdgcnspirv (gpu-generic) target needs for per-arch device libraries. Add DISABLE_TARGET_PROJECTS as a first-class primitive: projects listed under it are not declared/built for the target. Replace the ad-hoc therock_project_has_no_targets helper (which inspected the EXCLUDE list) with therock_project_disabled, keyed off the new DISABLE metadata. Split the amdgcnspirv project list: device libraries and their tests/providers/samples move to DISABLE (fully skipped); core runtime/tools stay in EXCLUDE (build host-only). composable_kernel stays in EXCLUDE because its _ck_enabled gate reuses the EXCLUDE-based filter shared with partial-support gfx shards. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
✅ All Checks Passed — Ready for Review
📖 Need help? See the Policy FAQ for details on every check and how to fix failures. |
|
🎉 All checks passed! This PR is ready for review. |
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.
Problem
EXCLUDE_TARGET_PROJECTSonly strips a gfx target from a project'sGPU_TARGETS— the project still builds (host-only or viaDEFAULT_GPU_TARGETSfallback). It cannot express "this project is meaningless for this target and must not build at all", which is exactly what the build-onlyamdgcnspirv(gpu-generic) target needs for per-arch device libraries. #6843 worked around this with an ad-hoctherock_project_has_no_targetshelper that inspected the EXCLUDE list.Change
Introduce
DISABLE_TARGET_PROJECTSas a first-class primitive:EXCLUDE_TARGET_PROJECTSGPU_TARGETS, project still buildsDISABLE_TARGET_PROJECTStherock_add_amdgpu_targetrecordsDISABLE_TARGET_PROJECTSintoTHEROCK_AMDGPU_PROJECT_TARGET_DISABLES_<name>.therock_project_disabledreplacestherock_project_has_no_targets(introduced in this PR chain, no external users); all gate call sites updated.amdgcnspirvproject list is split: device libraries and their tests/providers/samples → DISABLE (fully skipped); core runtime/tools → EXCLUDE (build host-only).composable_kernelstays in EXCLUDE because its_ck_enabledgate reuses the EXCLUDE-based filter shared with partial-support gfx shards.Behavior is preserved (the workaround and the primitive skip the same blocks under spirv), but the intent is now explicit and reusable for any future architecture-independent target.
Testing
python3 -m unittest build_tools.tests.cmake_amdgpu_targets_test— 11 pass (incl. new parser case for FAMILY terminating atDISABLE_TARGET_PROJECTS).cmake -Pharness exercisingtherock_add_amdgpu_target+therock_project_disabled— 6 semantic cases pass (device lib disabled under spirv; core-runtime not disabled; not disabled under gfx942; not disabled in mixed gfx942+spirv; not disabled in no-targets generic stage; unknown project not disabled).See
docs/10_disable_target_projects_plan.md.🤖 Generated with Claude Code