Skip to content
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.

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

refactor: replace old warp size api #705

Open
wants to merge 10 commits into
base: develop
Choose a base branch
from

Conversation

Naraenda
Copy link
Member

Required steps to get rid of __AMDGCN_WAVEFRONT_SIZE. This PR should be included in ROCm 6.5.

  • Using AMD nomenclature for wavefront sizes instead of warp sizes.
  • Opted to use namespace scopes, i.e. ::rocprim::arch::wavefront::min_size():
    • ::rocprim::wavefront::size() felt strange as it really isn't a top level API (i.e. algorithm)
    • ::rocprim::arch::wavefront_size() was also a possibility, but we might add more wavefront related functions. So scoping it made more sense.
    • ::rocprim::arch was more appropriate so we can add more architecture related constants such as number of LDS banks, etc.
  • I basically did a search and replace. min_size() and max_size() are the same on device for now.

@Naraenda Naraenda self-assigned this Mar 17, 2025
@Naraenda
Copy link
Member Author

Still in draft as we're doing an internal review in parallel. @stanleytsang-amd @RobsonRLemos feel to already run this through the tests to speed things along as I don't think much will change.

I'll probably open up a pair of PRs to hipCUB and rocThrust to fix the compiler warnings from the deprecations tomorrow.

@Naraenda Naraenda changed the title Deprecate old warp size refactor: replace old warp size api Mar 17, 2025
@Naraenda Naraenda force-pushed the deprecate-old-warp-size branch from 8046d96 to 76f60f1 Compare March 19, 2025 11:35
@Naraenda Naraenda changed the base branch from develop to release-staging/rocm-rel-6.5 March 19, 2025 12:04
@Naraenda Naraenda changed the base branch from release-staging/rocm-rel-6.5 to develop March 19, 2025 12:04
@Naraenda Naraenda force-pushed the deprecate-old-warp-size branch from 76f60f1 to 7765574 Compare March 19, 2025 13:25
@Naraenda Naraenda marked this pull request as ready for review March 19, 2025 13:26
@Naraenda
Copy link
Member Author

While applying the migration to the new API on hipCUB, I've found an unrollable loop which is fixed in 7765574.

@stanleytsang-amd @RobsonRLemos Apologies for the delay, there were a few oversights that took a bit longer to resolve.

…/warp algo types on host
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.

None yet

2 participants