Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
carlushuang wants to merge 1 commit into
base: release/v0.1.17
Choose a base branch
from
Open

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
7 changes: 4 additions & 3 deletions csrc/include/opus/opus.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1627,8 +1627,8 @@ OPUS_D unsigned int lane_id() {
else return __builtin_amdgcn_mbcnt_hi(-1, __builtin_amdgcn_mbcnt_lo(-1, 0));
}

//gfx1250 only feature
#if defined(__gfx1250__) || !defined(__HIP_DEVICE_COMPILE__)
//gfx1250 only feature. Named-barrier / cluster-sync builtins + __amdgpu_named_workgroup_barrier_t only exist on clang>=22 (ROCm>=7.2); gate the host pass too so clang-20 (ROCm 7.1) CI does not parse them.
#if (defined(__gfx1250__) || !defined(__HIP_DEVICE_COMPILE__)) && (__clang_major__ >= 22)
OPUS_D u32_t waveid_in_workgroup() { u32_t wave_id; asm volatile("s_bfe_u32 %0, ttmp8, 0x50019" : "=s"(wave_id)); return wave_id; }

//Named Barrier define
Expand Down Expand Up @@ -2328,7 +2328,8 @@ template<typename T_> OPUS_D decltype(auto) make_smem(T_* ptr) { return smem<T_>

/////////////////////////////////////////////////////////////////////////////////////////////////////////
// tdm (gfx1250): tdm_desc = stateless D# (sg0..sg3 raw dwords); tdm_window = stateful tile window with make() + move(d0..d4, lds) + load_to_lds<cpol>(); cpol[6:0] = | rsvd | NV | scope[2] | th[3] |.
#if defined(__gfx1250__) || !defined(__HIP_DEVICE_COMPILE__)
// __builtin_amdgcn_tensor_load_to_lds only exists on clang>=22 (ROCm>=7.2); gate the host pass too so clang-20 (ROCm 7.1) CI does not parse it.
#if (defined(__gfx1250__) || !defined(__HIP_DEVICE_COMPILE__)) && (__clang_major__ >= 22)

enum class tdm_load_th : u8_t { rt=0, nt=1, ht=2, bypass=3, nt_rt=4, rt_nt=5, nt_ht=6 }; // bypass = LU (last-use)
enum class tdm_scope : u8_t { cu=0, se=1, dev=2, sys=3 };
Expand Down
7 changes: 4 additions & 3 deletions csrc/opus_gemm/include/gfx1250/opus_gemm_traits_a16w16_gfx1250.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -261,9 +261,10 @@ struct opus_cluster_tdm_splitk_ws_traits_gfx1250 {
static constexpr int kSchedWmmaMask = 0x008; // WMMA group
static constexpr int kSchedWmmaCount = kExpM * kExpN * kExpKHalf;

#if defined(__gfx1250__) || !defined(__HIP_DEVICE_COMPILE__)
// tdm_window types live only where tdm_window is available (gfx1250
// device pass + host pass). Non-gfx1250 device passes never reference them.
#if (defined(__gfx1250__) || !defined(__HIP_DEVICE_COMPILE__)) && (__clang_major__ >= 22)
// tdm_window types live only where tdm_window is available (gfx1250 device
// pass + host pass, clang>=22). Non-gfx1250 device passes and clang-20
// (ROCm 7.1) CI never reference them -- opus::tdm_window is gated identically.
using WindowA = opus::tdm_window<DataA, kTdmK, kARows, 0, 0, 0,
1, 0, 0, 0, 1, 0, 0, 0, 0,
kLdsPadEn, kPadInterval, kPadAmount, opus::seq<>>;
Expand Down