Bugfix/gh 1901 scan algo by oleksandr-pavlyk · Pull Request #1902 · IntelPython/dpctl · GitHub
Skip to content

Bugfix/gh 1901 scan algo #1902

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 4 commits into from
Nov 21, 2024
Merged

Bugfix/gh 1901 scan algo #1902

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
ee7d6fb
Stylistic changes in accumulators.hpp
oleksandr-pavlyk Nov 20, 2024
d79831e
Fix a lapse in logic for computing upper bound of iteration
oleksandr-pavlyk Nov 20, 2024
abb1572
Add changelog entry for fixing of gh-1901
oleksandr-pavlyk Nov 20, 2024
d133a11
Add a test based on example from gh-1901
oleksandr-pavlyk Nov 20, 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
1 change: 1 addition & 0 deletions CHANGELOG.md
View file
Open in desktop
22 changes: 12 additions & 10 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 @@ -215,16 +215,17 @@ inclusive_scan_base_step(sycl::queue &exec_q,
const size_t gid = it.get_global_id(0);
const size_t lid = it.get_local_id(0);

const size_t iter_gid = gid / (acc_groups * wg_size);
const size_t chunk_gid = gid - (iter_gid * acc_groups * wg_size);
const size_t reduce_chunks = acc_groups * wg_size;
const size_t iter_gid = gid / reduce_chunks;
const size_t chunk_gid = gid - (iter_gid * reduce_chunks);

std::array<outputT, n_wi> local_iscan;

size_t i = chunk_gid * n_wi;
const size_t i = chunk_gid * n_wi;
const auto &iter_offsets = iter_indexer(iter_gid);
const auto &inp_iter_offset = iter_offsets.get_first_offset();
const auto &out_iter_offset = iter_offsets.get_second_offset();

std::array<outputT, n_wi> local_iscan;

#pragma unroll
for (nwiT m_wi = 0; m_wi < n_wi; ++m_wi) {
const size_t i_m_wi = i + m_wi;
Expand Down Expand Up @@ -279,8 +280,9 @@ inclusive_scan_base_step(sycl::queue &exec_q,
local_iscan[m_wi] = scan_op(local_iscan[m_wi], addand);
}

const nwiT m_max =
std::min<nwiT>(n_wi, std::max(i, acc_nelems) - i);
const size_t start = std::min(i, acc_nelems);
const size_t end = std::min(i + n_wi, acc_nelems);
const nwiT m_max = static_cast<nwiT>(end - start);
for (nwiT m_wi = 0; m_wi < m_max; ++m_wi) {
output[out_iter_offset + out_indexer(i + m_wi)] =
local_iscan[m_wi];
Expand Down Expand Up @@ -324,7 +326,7 @@ sycl::event inclusive_scan_iter_1d(sycl::queue &exec_q,
std::vector<sycl::event> &host_tasks,
const std::vector<sycl::event> &depends = {})
{
ScanOpT scan_op = ScanOpT();
ScanOpT scan_op{};
constexpr outputT identity = su_ns::Identity<ScanOpT, outputT>::value;

constexpr size_t _iter_nelems = 1;
Expand Down Expand Up @@ -352,9 +354,9 @@ sycl::event inclusive_scan_iter_1d(sycl::queue &exec_q,
size_t n_groups_ = n_groups;
size_t temp_size = 0;
while (n_groups_ > 1) {
const auto this_size = (n_groups_ - 1);
const size_t this_size = (n_groups_ - 1);
temp_size += this_size;
n_groups_ = ceiling_quotient<size_t>(this_size, chunk_size);
n_groups_ = ceiling_quotient(this_size, chunk_size);
}

// allocate
Expand Down
11 changes: 11 additions & 0 deletions dpctl/tests/test_tensor_accumulation.py
View file
Open in desktop