Skip to content

Implements top_k in dpctl.tensor #1921

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
ndgrigorian merged 30 commits into from
Jan 5, 2025
Merged

Implements top_k in dpctl.tensor #1921

Changes from 1 commit
Commits
Show all changes
30 commits
Select commit Hold shift + click to select a range
0b353bf
Implements `top_k`
ndgrigorian Dec 4, 2024
f809573
Add `top_k` to doc sources
ndgrigorian Dec 5, 2024
b59e59d
Use `std::size_t` for `k` instead of `py::ssize_t`
ndgrigorian Dec 5, 2024
b4d7ba4
Add implementation of `top_k` using radix sort
ndgrigorian Dec 5, 2024
c37f1de
Reduce code duplication by using std::optional in py_topk directly
ndgrigorian Dec 11, 2024
1f0e9fd
Clean up top_k docstring
ndgrigorian Dec 11, 2024
fa74e5e
Move top_k into _sorting.py
ndgrigorian Dec 12, 2024
cd83881
top_k raises when k > 1 and input is a scalar
ndgrigorian Dec 13, 2024
70c2ef0
Fix bug in top_k partial merge sort implementation
ndgrigorian Dec 13, 2024
65f14be
Add test file for top_k functionality
oleksandr-pavlyk Dec 14, 2024
275827a
Add sort_utils.hpp with iota_impl
oleksandr-pavlyk Dec 20, 2024
cd1243f
Simplify write-out kernels in topk implementation (avoid recomputing …
oleksandr-pavlyk Dec 22, 2024
823c201
Add explicit read into registers
oleksandr-pavlyk Dec 22, 2024
3a4d8ab
Use unique_ptr as temporary owner of USM allocation
oleksandr-pavlyk Dec 20, 2024
b15e302
Factor out write-out kernel into separate detail function
oleksandr-pavlyk Dec 25, 2024
4ef615e
Typo: smart_malloc_jost -> smart_malloc_host
oleksandr-pavlyk Dec 25, 2024
2608aea
Reimplemented map_back_impl to process few elements per work-item
oleksandr-pavlyk Dec 25, 2024
15da303
Replace use of std::log2(size_t_value)
oleksandr-pavlyk Dec 25, 2024
fb9bc43
Remove dead code disable via preprocessor conditional
oleksandr-pavlyk Dec 27, 2024
6cf36b1
Add information displayed on failure, renamed variables
oleksandr-pavlyk Dec 25, 2024
4ec252e
Add comment to explain ceil_log2 algo
oleksandr-pavlyk Dec 27, 2024
1436522
Add static assers to async_smart_free
oleksandr-pavlyk Dec 27, 2024
f0e3245
Skip top_k test known to fail for some CPU architectures
ndgrigorian Dec 28, 2024
8f4dcdd
Fixed blunder in work-item id to data_id computation
oleksandr-pavlyk Dec 31, 2024
b6806ba
Replace map_back_impl in sort_utils
oleksandr-pavlyk Dec 31, 2024
b47d9f3
Use get_global_linear_id instead of get_global_id and rely on implici…
oleksandr-pavlyk Jan 2, 2025
0493485
Counters in one-workgroup kernel to use uint16_t from uint32_t
oleksandr-pavlyk Jan 2, 2025
5125e11
Apply work-around for failing tests with CPU device and short sub-groups
oleksandr-pavlyk Jan 4, 2025
505b64c
Remove skipping of tests for i1/i2 dtypes since work-around
oleksandr-pavlyk Dec 28, 2024
8c6abf5
Add entry to changelog for top_k
ndgrigorian Jan 5, 2025
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
Fixed blunder in work-item id to data_id computation 
gid-lane_id is already a multiple of sg_size.
  • Loading branch information
@oleksandr-pavlyk @ndgrigorian
oleksandr-pavlyk authored and ndgrigorian committed Jan 5, 2025
commit 8f4dcdd05fb1e5118da91810aee08b7cb700efdf
3 changes: 1 addition & 2 deletions dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -101,8 +101,7 @@ sycl::event write_out_impl(sycl::queue &exec_q,
const std::uint32_t lane_id = sg.get_local_id()[0];
const std::uint32_t sg_size = sg.get_max_local_range()[0];

const std::size_t start_id =
(gid - lane_id) * sg_size * n_wi + lane_id;
const std::size_t start_id = (gid - lane_id) * n_wi + lane_id;

#pragma unroll
for (std::uint32_t i = 0; i < n_wi; ++i) {
Expand Down