Skip to content

[opus] backport #4056: gate TDM/named-barrier on clang>=22 for ROCm 7.0/7.1 (release/v0.1.17)#4078

Open
carlushuang wants to merge 1 commit into
ROCm:release/v0.1.17from
carlushuang:carhuang/backport-4056-tdm-guard-v0.1.17
Open

[opus] backport #4056: gate TDM/named-barrier on clang>=22 for ROCm 7.0/7.1 (release/v0.1.17)#4078
carlushuang wants to merge 1 commit into
ROCm:release/v0.1.17from
carlushuang:carhuang/backport-4056-tdm-guard-v0.1.17

Conversation

@carlushuang

Copy link
Copy Markdown
Collaborator

Backport of #4056 into release/v0.1.17 to fix #4077.

release/v0.1.17 HEAD (68b5f7be) fails to build on ROCm 7.0/7.1 (clang-20) for gfx942;gfx950 because the gfx1250-only TDM section (tdm_desc/tdm_window/__builtin_amdgcn_tensor_load_to_lds) and the named-barrier / cluster-sync block (__amdgpu_named_workgroup_barrier_t, s_barrier_init) in opus.hpp had their host-pass branch (!defined(__HIP_DEVICE_COMPILE__)) compiled unconditionally. Those builtins only exist on clang>=22 (ROCm>=7.2), so any .cu pulling in opus.hpp (e.g. module_activation) failed with use of undeclared identifier '__builtin_amdgcn_tensor_load_to_lds'.

This cherry-picks the fix already merged to main (#4056), adding && (__clang_major__ >= 22) to both guards in opus.hpp and the matching WindowA/WindowB aliases in opus_gemm_traits_a16w16_gfx1250.cuh, so the whole block compiles out on clang-20 while remaining active on ROCm 7.2.

Cherry-pick applied cleanly (no conflicts). Verified guard counts after cherry-pick: opus.hpp = 2, traits header = 1 (matching main).

Fixes #4077.

…builds (ROCm#4056)

The gfx1250 TDM section (tdm_desc/tdm_window/__builtin_amdgcn_tensor_load_to_lds)
and the named-barrier / cluster-sync block (__amdgpu_named_workgroup_barrier_t,
s_barrier_init) only exist on clang>=22 (ROCm>=7.2). Their host-pass branch
(!defined(__HIP_DEVICE_COMPILE__)) was compiled unconditionally, so clang-20
(ROCm 7.1) CI failed with "use of undeclared identifier
'__builtin_amdgcn_tensor_load_to_lds'". Add && (__clang_major__ >= 22) to those
guards (opus.hpp x2) and the matching tdm_window WindowA/WindowB aliases in the
traits header so the whole block compiles out on clang-20.

(cherry picked from commit e2ff4b3)
@zufayu zufayu requested a review from yzhou103 July 6, 2026 02:01
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