Skip to content

Improve base step of scan algorithm #1942

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
oleksandr-pavlyk merged 6 commits into from
Dec 27, 2024
Merged

Improve base step of scan algorithm #1942

Changes from 1 commit
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
18c5ed7
Add implementation of base scan using striped load/store pattern
oleksandr-pavlyk Dec 18, 2024
ffd2609
Rename variable in inclusive_scan_base_step_blocked, add const qualifier
oleksandr-pavlyk Dec 19, 2024
101514f
Implement custom scan group function for generic binary operator
oleksandr-pavlyk Dec 19, 2024
ce4cf5d
Fixed hang on CUDA GPUs in custom scanning code
oleksandr-pavlyk Dec 20, 2024
0d59c8f
For HIP devices use smaller wg-size parameter
oleksandr-pavlyk Dec 22, 2024
93eba50
Changed type of wg_size to uint32_t everywhere
oleksandr-pavlyk Dec 26, 2024
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
Prev Previous commit
Next Next commit
For HIP devices use smaller wg-size parameter 
single_step_scan_striped does not produce correct results for
wg_size > 64, and tests fail.
  • Loading branch information
@oleksandr-pavlyk
oleksandr-pavlyk committed Dec 22, 2024
commit 0d59c8f51e3bd5ccbeb64ac1b132fc78da42394f
20 changes: 16 additions & 4 deletions dpctl/tensor/libtensor/include/kernels/accumulators.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -777,7 +777,10 @@ accumulate_1d_contig_impl(sycl::queue &q,
}
else {
constexpr nwiT n_wi_for_gpu = 4;
const std::size_t wg_size = 256;
// base_scan_striped algorithm does not execute correctly
// on HIP device with wg_size > 64
const std::size_t wg_size =
(q.get_backend() == sycl::backend::ext_oneapi_hip) ? 64 : 256;
comp_ev = inclusive_scan_iter_1d<srcT, dstT, n_wi_for_gpu, NoOpIndexerT,
transformerT, AccumulateOpT,
include_initial>(
Expand Down Expand Up @@ -1181,7 +1184,10 @@ accumulate_strided_impl(sycl::queue &q,
}
else {
constexpr nwiT n_wi_for_gpu = 4;
const std::size_t wg_size = 256;
// base_scan_striped algorithm does not execute correctly
// on HIP device with wg_size > 64
const std::size_t wg_size =
(q.get_backend() == sycl::backend::ext_oneapi_hip) ? 64 : 256;
comp_ev =
inclusive_scan_iter<srcT, dstT, n_wi_for_gpu, InpIndexerT,
OutIndexerT, InpIndexerT, OutIndexerT,
Expand Down Expand Up @@ -1235,7 +1241,10 @@ std::size_t cumsum_val_contig_impl(sycl::queue &q,
}
else {
constexpr nwiT n_wi_for_gpu = 4;
const std::size_t wg_size = 256;
// base_scan_striped algorithm does not execute correctly
// on HIP device with wg_size > 64
const std::size_t wg_size =
(q.get_backend() == sycl::backend::ext_oneapi_hip) ? 64 : 256;
comp_ev = inclusive_scan_iter_1d<maskT, cumsumT, n_wi_for_gpu,
NoOpIndexerT, transformerT,
AccumulateOpT, include_initial>(
Expand Down Expand Up @@ -1346,7 +1355,10 @@ cumsum_val_strided_impl(sycl::queue &q,
}
else {
constexpr nwiT n_wi_for_gpu = 4;
const std::size_t wg_size = 256;
// base_scan_striped algorithm does not execute correctly
// on HIP device with wg_size > 64
const std::size_t wg_size =
(q.get_backend() == sycl::backend::ext_oneapi_hip) ? 64 : 256;
comp_ev = inclusive_scan_iter_1d<maskT, cumsumT, n_wi_for_gpu,
StridedIndexerT, transformerT,
AccumulateOpT, include_initial>(
Expand Down