Skip to content

Add class type support for __builtin_amdgcn_readfirstlane()#711

Merged
asroy merged 18 commits into
developfrom
feature/support-readfirstlane-for-object-types
May 31, 2023
Merged

Add class type support for __builtin_amdgcn_readfirstlane()#711
asroy merged 18 commits into
developfrom
feature/support-readfirstlane-for-object-types

Conversation

@poyenc
Copy link
Copy Markdown
Contributor

@poyenc poyenc commented May 16, 2023

This PR added ck::amd_wave_read_first_lane() overloaded functions to unify call syntax of __builtin_amdgcn_readfirstlane().
The ck::amd_wave_read_first_lane() supports std::int32_t & trivially copyable object arguments.

@poyenc poyenc self-assigned this May 16, 2023
@poyenc poyenc requested a review from asroy May 16, 2023 19:40
@poyenc poyenc added the WIP label May 18, 2023
@poyenc
Copy link
Copy Markdown
Contributor Author

poyenc commented May 18, 2023

It seems the offset calculation in static_for<> is wrong. Need to figure out right boundary.

@poyenc poyenc removed the WIP label May 18, 2023
Comment thread include/ck/utility/readfirstlane.hpp Outdated
Comment thread include/ck/utility/readfirstlane.hpp Outdated
Comment thread include/ck/utility/readfirstlane.hpp Outdated
@poyenc
Copy link
Copy Markdown
Contributor Author

poyenc commented May 29, 2023

I mode following changes

  • Replace static_for<> by ordinary for loop
  • Rename type trait from get_signed_int<> to get_unsigned_int<>
  • Remove std:: qualifiers from standard integral types
  • Replace size_t by unsigned to represent object size (non-negative)

@zjing14 zjing14 requested a review from asroy May 30, 2023 12:17
Comment thread include/ck/utility/readfirstlane.hpp Outdated
template <
typename Object,
typename = std::enable_if_t<std::is_class_v<Object> && std::is_trivially_copyable_v<Object>>>
__device__ auto readfirstlane(const Object& obj)
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

This is a AMD specific function, not standard HIP. I think better rename to "amd_wave_read_first_lane", similar to "amd_buffer_load"

@asroy asroy self-requested a review May 30, 2023 16:54
Copy link
Copy Markdown
Contributor

@asroy asroy left a comment

Choose a reason for hiding this comment

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

Please rename function from to "amd_xxx", maybe "amd_wave_read_first_lane". Otherwise LGTM

@poyenc
Copy link
Copy Markdown
Contributor Author

poyenc commented May 31, 2023

Please rename function from to "amd_xxx", maybe "amd_wave_read_first_lane". Otherwise LGTM

Done. I also renamed the header file.

@asroy asroy self-requested a review May 31, 2023 15:24
@asroy asroy merged commit 582e31e into develop May 31, 2023
@illsilin illsilin deleted the feature/support-readfirstlane-for-object-types branch December 7, 2023 18:56
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