Skip to content

Sync Barrier for persistent kernel #2977

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

Merged
ThomasNing merged 1 commit into from
Oct 7, 2025
Merged

Sync Barrier for persistent kernel #2977

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
1 change: 1 addition & 0 deletions include/ck_tile/ops/gemm/kernel/universal_gemm_kernel.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1134,6 +1134,7 @@ struct UniversalGemmKernel

while(block_id < num_work)
{
s_waitcnt_barrier();
Copy link
Copy Markdown
Collaborator

@aosewski aosewski Oct 6, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

That's interesting, cause it has been working for all gfx9 gpus. Shouldn't we conditionally add this for other cards?

Copy link
Copy Markdown
Contributor Author

@ThomasNing ThomasNing Oct 7, 2025

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It fails intermittently on the gfx11 series. As it is blocking with other PRs, I will merge it first and add the device condition on the next PR.

// Get the tile index for this block
const auto tile_idx = amd_wave_read_first_lane(block_id % num_tiles);
const auto [iM, iN] = TilePartitioner{kargs.M, kargs.N}.GetOutputTileIndex(tile_idx);
Expand Down
Loading