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.

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

Jump to bottom

Fix hip scan accumulation types #1634

Closed
wants to merge 4 commits into from
Closed

Fix hip scan accumulation types #1634

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
eace150
Change init type in exclusive_inplace
MrBurmark Apr 27, 2024
82cb3b6
Change init type in exclusive
MrBurmark Apr 27, 2024
3060a9f
Change AccType in inclusive
MrBurmark Apr 27, 2024
f920936
fixup commas used in hipErrchk macro
MrBurmark Apr 27, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
51 changes: 35 additions & 16 deletions include/RAJA/policy/hip/scan.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -135,10 +135,13 @@ exclusive_inplace(
InputIter begin,
InputIter end,
Function binary_op,
T init)
T init_in)
{
hipStream_t stream = hip_res.get_stream();

using AccType = typename std::iterator_traits<InputIter>::value_type;
auto init = static_cast<AccType>(init_in);

int len = std::distance(begin, end);
// Determine temporary device storage requirements
void* d_temp_storage = nullptr;
Expand Down Expand Up @@ -217,18 +220,27 @@ inclusive(
{
hipStream_t stream = hip_res.get_stream();

#if defined(__HIPCC__)
using Config = rocprim::default_config;
using AccType = typename std::iterator_traits<OutputIter>::value_type;
#endif

int len = std::distance(begin, end);
// Determine temporary device storage requirements
void* d_temp_storage = nullptr;
size_t temp_storage_bytes = 0;
#if defined(__HIPCC__)
hipErrchk(::rocprim::inclusive_scan(d_temp_storage,
temp_storage_bytes,
begin,
out,
len,
binary_op,
stream));
hipErrchk((::rocprim::inclusive_scan<Config,
InputIter,
OutputIter,
Function,
AccType>(d_temp_storage,
Copy link
Member

Choose a reason for hiding this comment

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

I think the compilation failure is because the ROCm template you are calling here only has 4 template params and you have five. Specifically, AccType does not appear in the ROCm implementation.

Copy link
Member Author

@MrBurmark MrBurmark Apr 29, 2024 •
edited

Choose a reason for hiding this comment

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

Ya, I just realized that.
It looks like all the released versions of rocprim don't have the 5th AccType template argument as of rocm 6.1.0. Though this template argument does exist on rocprim develop.

Copy link
Member

Choose a reason for hiding this comment

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

It's also not in rocm 5.7.x, which is what we are currently testing.

temp_storage_bytes,
begin,
out,
len,
binary_op,
stream)));
#elif defined(__CUDACC__)
hipErrchk(::cub::DeviceScan::InclusiveScan(d_temp_storage,
temp_storage_bytes,
Expand All @@ -244,13 +256,17 @@ inclusive(
temp_storage_bytes);
// Run
#if defined(__HIPCC__)
hipErrchk(::rocprim::inclusive_scan(d_temp_storage,
temp_storage_bytes,
begin,
out,
len,
binary_op,
stream));
hipErrchk((::rocprim::inclusive_scan<Config,
InputIter,
OutputIter,
Function,
AccType>(d_temp_storage,
temp_storage_bytes,
begin,
out,
len,
binary_op,
stream)));
#elif defined(__CUDACC__)
hipErrchk(::cub::DeviceScan::InclusiveScan(d_temp_storage,
temp_storage_bytes,
Expand Down Expand Up @@ -289,10 +305,13 @@ exclusive(
InputIter end,
OutputIter out,
Function binary_op,
T init)
T init_in)
{
hipStream_t stream = hip_res.get_stream();

using AccType = typename std::iterator_traits<OutputIter>::value_type;
auto init = static_cast<AccType>(init_in);

int len = std::distance(begin, end);
// Determine temporary device storage requirements
void* d_temp_storage = nullptr;
Expand Down