Implements `top_k` in dpctl.tensor #1921

ndgrigorian · 2024-12-04T10:38:49Z

This PR implements the functions top_k as per proposal in array API spec.

Radix and merge sorting are used, and modified merge-sort kernels are introduced which sort the array in chunks and write out to a temporary the k largest or smallest values.

Have you provided a meaningful PR description?
Have you added a test, reproducer or referred to an issue with a reproducer?
Have you tested your changes locally for CPU and GPU devices?
Have you made sure that new changes do not introduce compiler warnings?
Have you checked performance impact of proposed changes?
Have you added documentation for your changes, if necessary?
Have you added your changes to the changelog?
If this PR is a work in progress, are you opening the PR as a draft?

github-actions · 2024-12-04T11:16:50Z

Deleted rendered PR docs from intelpython.github.com/dpctl, latest should be updated shortly. 🤞

github-actions · 2024-12-04T11:26:19Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310hdf72452_296 ran successfully.
Passed: 895
Failed: 1
Skipped: 118

coveralls · 2024-12-04T11:27:46Z

coverage: 87.716% (+0.06%) from 87.659%
when pulling 8c6abf5 on feature/topk
into f7cb1b1 on master.

oleksandr-pavlyk · 2024-12-04T21:10:49Z

@ndgrigorian Please add top_k to the docs/doc_sources/api_reference/dpctl/tensor.sorting_functions.rst:

diff --git a/docs/doc_sources/api_reference/dpctl/tensor.sorting_functions.rst b/docs/doc_sources/api_reference/dpctl/tensor.sorting_functions.rst
index ae1605d988..ef20f4654c 100644
--- a/docs/doc_sources/api_reference/dpctl/tensor.sorting_functions.rst
+++ b/docs/doc_sources/api_reference/dpctl/tensor.sorting_functions.rst
@@ -10,3 +10,4 @@ Sorting functions

    argsort
    sort
+   top_k

github-actions · 2024-12-05T01:00:55Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310hdf72452_295 ran successfully.
Passed: 894
Failed: 2
Skipped: 118

github-actions · 2024-12-06T00:08:39Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310hdf72452_297 ran successfully.
Passed: 895
Failed: 1
Skipped: 118

github-actions · 2024-12-11T19:31:47Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_326 ran successfully.
Passed: 895
Failed: 1
Skipped: 118

github-actions · 2024-12-11T22:19:48Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_327 ran successfully.
Passed: 894
Failed: 2
Skipped: 118

github-actions · 2024-12-12T10:09:12Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_331 ran successfully.
Passed: 894
Failed: 2
Skipped: 118

github-actions · 2024-12-13T01:42:26Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_331 ran successfully.
Passed: 895
Failed: 1
Skipped: 118

github-actions · 2024-12-18T22:28:52Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_331 ran successfully.
Passed: 894
Failed: 2
Skipped: 118

github-actions · 2024-12-27T19:40:03Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_388 ran successfully.
Passed: 893
Failed: 3
Skipped: 118

github-actions · 2024-12-28T01:29:31Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_391 ran successfully.
Passed: 894
Failed: 2
Skipped: 118

github-actions · 2024-12-28T01:43:20Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_386 ran successfully.
Passed: 894
Failed: 2
Skipped: 118

github-actions · 2024-12-28T01:56:53Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_385 ran successfully.
Passed: 895
Failed: 1
Skipped: 118

github-actions · 2024-12-28T01:59:08Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_386 ran successfully.
Passed: 895
Failed: 1
Skipped: 118

github-actions · 2024-12-28T02:52:06Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_387 ran successfully.
Passed: 894
Failed: 2
Skipped: 118

ndgrigorian · 2024-12-28T05:27:43Z

@oleksandr-pavlyk
I've added a skip for both the smallest and largest tests with int8

oleksandr-pavlyk · 2024-12-28T14:30:33Z

Good to see the CI green again!

I was suggesting to only skip it in the test command used in the workflow so that we can still provide a reproducer to the CPU team.

I was thinking we could add a file of tests to skip, and pass it as argument to pytest.

I think this is the approach taken by dpnp.

…gid)

Until it is passed over to the host function, and unique_ptr's ownership is released. Also reduced allocation sizes, where too much was being allocated. Introduce smart_malloc_device, etc. The smart_malloc_device<T>(count, q) makes USM allocation and returns a unique_ptr<T, USMDeleter> which owns the allocation. The function throws an exception (std::runtime_error) if USM allocation is not successful. Introduce async_smart_free. This function intends to replace use of host_task submissions to manage USM temporary deallocations. The usage is as follows: ``` // returns unique_ptr auto alloc_owner = smart_malloc_device<T>(count, q); // get raw pointer for use in kernels T *data = alloc_owner.get(); [..SNIP..] // submit host_task that releases the unique_ptr // after the host task was successfully submitted // and ownership of USM allocation is transfered to // the said host task sycl::event ht_ev = async_smart_free(q, dependent_events, alloc_owner); [...SNIP...] ```

Replaced three duplicates of the same kernel with calls to this function.

Factored out map_back_impl projects indexing from flat index to a row-wise index. Removed dead code excluded by preprocessor conditional.

Replaced it with hand-written implementation of ceil_log2(n), such that n <= (dectype(n){1} << ceil_log2(n)) is true for all positive values of `n` in the range.

Add check of computed against expected indices

One asserts that at least one unique pointer is specified. Another that specified arguments are unique pointers with USMDeleter.

gid-lane_id is already a multiple of sg_size.

Change kernel to process few data elements in the work-item.

…t conversion

Counters can not exceed uint16_t max, because the kernel assumes that the number of elements to sort fits into uint16_t. The change reduces the kernel SLM footprint. Also, remove use of std::move, uint16_t->std::uint16_t, etc Replace size_t->std::size_t, uint32_t->std::uint32_t Use `if constexpr` in order-preservign-cast for better readability.

The team developing OpenCL:CPU device runtime and compiler was notified. See CMPLRLLVM-64592 Once fixed, the work-around should be removed.

was applied in C++. Add tests for 2d input arrays, for axis=0 and axis=1 Add a test for non-contiguous input, 0d input, validation 100% coverage of top_k function implementation achieved

ndgrigorian · 2025-01-05T01:18:37Z

I'd suggest rebasing the branch on top of the targeted base branch to remove two cherry-picked commits fixing the workflow for building with nightly DPC++ bundle

Done, I will add the addition of topk to the changelog

github-actions · 2025-01-05T02:00:47Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_395 ran successfully.
Passed: 892
Failed: 4
Skipped: 118

github-actions · 2025-01-05T21:54:27Z

Array API standard conformance tests for dpctl=0.19.0dev0=py310h93fe807_396 ran successfully.
Passed: 893
Failed: 3
Skipped: 118

ndgrigorian requested a review from AlexanderKalistratov December 4, 2024 10:38

ndgrigorian force-pushed the feature/topk branch from 79b97d9 to 882c70d Compare December 5, 2024 00:16

ndgrigorian force-pushed the feature/topk branch from a56e21c to 26718f3 Compare December 11, 2024 18:46

ndgrigorian marked this pull request as ready for review December 12, 2024 09:24

ndgrigorian requested a review from oleksandr-pavlyk as a code owner December 12, 2024 09:24

ndgrigorian force-pushed the feature/topk branch 2 times, most recently from 8bcb100 to 8f38b80 Compare December 13, 2024 01:00

ndgrigorian force-pushed the feature/topk branch 4 times, most recently from 84d1388 to 809cb70 Compare December 28, 2024 01:13

oleksandr-pavlyk and others added 18 commits January 4, 2025 17:16

Simplify write-out kernels in topk implementation (avoid recomputing …

cd1243f

…gid)

Add explicit read into registers

823c201

Factor out write-out kernel into separate detail function

b15e302

Replaced three duplicates of the same kernel with calls to this function.

Typo: smart_malloc_jost -> smart_malloc_host

4ef615e

Reimplemented map_back_impl to process few elements per work-item

2608aea

Factored out map_back_impl projects indexing from flat index to a row-wise index. Removed dead code excluded by preprocessor conditional.

Replace use of std::log2(size_t_value)

15da303

Replaced it with hand-written implementation of ceil_log2(n), such that n <= (dectype(n){1} << ceil_log2(n)) is true for all positive values of `n` in the range.

Remove dead code disable via preprocessor conditional

fb9bc43

Add information displayed on failure, renamed variables

6cf36b1

Add check of computed against expected indices

Add comment to explain ceil_log2 algo

4ec252e

Add static assers to async_smart_free

1436522

One asserts that at least one unique pointer is specified. Another that specified arguments are unique pointers with USMDeleter.

Skip top_k test known to fail for some CPU architectures

f0e3245

Fixed blunder in work-item id to data_id computation

8f4dcdd

gid-lane_id is already a multiple of sg_size.

Replace map_back_impl in sort_utils

b6806ba

Change kernel to process few data elements in the work-item.

Use get_global_linear_id instead of get_global_id and rely on implici…

b47d9f3

…t conversion

Apply work-around for failing tests with CPU device and short sub-groups

5125e11

The team developing OpenCL:CPU device runtime and compiler was notified. See CMPLRLLVM-64592 Once fixed, the work-around should be removed.

Remove skipping of tests for i1/i2 dtypes since work-around

505b64c

was applied in C++. Add tests for 2d input arrays, for axis=0 and axis=1 Add a test for non-contiguous input, 0d input, validation 100% coverage of top_k function implementation achieved

ndgrigorian force-pushed the feature/topk branch from 2865f27 to 505b64c Compare January 5, 2025 01:16

oleksandr-pavlyk previously approved these changes Jan 5, 2025

View reviewed changes

ndgrigorian changed the title ~~Implements top_k functions in dpctl.tensor~~ Implements top_k in dpctl.tensor Jan 5, 2025

Add entry to changelog for top_k

8c6abf5

ndgrigorian dismissed oleksandr-pavlyk’s stale review via 8c6abf5 January 5, 2025 21:08

ndgrigorian requested a review from oleksandr-pavlyk January 5, 2025 21:08

oleksandr-pavlyk approved these changes Jan 5, 2025

View reviewed changes

ndgrigorian merged commit 3a1a7c5 into master Jan 5, 2025
58 of 61 checks passed

ndgrigorian deleted the feature/topk branch January 5, 2025 23:06

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Implements `top_k` in dpctl.tensor #1921

Implements `top_k` in dpctl.tensor #1921

ndgrigorian commented Dec 4, 2024 •

edited

Loading

github-actions bot commented Dec 4, 2024 •

edited

Loading

github-actions bot commented Dec 4, 2024

coveralls commented Dec 4, 2024 •

edited

Loading

oleksandr-pavlyk commented Dec 4, 2024

github-actions bot commented Dec 5, 2024

github-actions bot commented Dec 6, 2024

github-actions bot commented Dec 11, 2024

github-actions bot commented Dec 11, 2024

github-actions bot commented Dec 12, 2024

github-actions bot commented Dec 13, 2024

github-actions bot commented Dec 18, 2024

github-actions bot commented Dec 27, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

ndgrigorian commented Dec 28, 2024

oleksandr-pavlyk commented Dec 28, 2024

ndgrigorian commented Jan 5, 2025

github-actions bot commented Jan 5, 2025

github-actions bot commented Jan 5, 2025

Implements top_k in dpctl.tensor #1921

Implements top_k in dpctl.tensor #1921

Conversation

ndgrigorian commented Dec 4, 2024 • edited Loading

github-actions bot commented Dec 4, 2024 • edited Loading

github-actions bot commented Dec 4, 2024

coveralls commented Dec 4, 2024 • edited Loading

oleksandr-pavlyk commented Dec 4, 2024

github-actions bot commented Dec 5, 2024

github-actions bot commented Dec 6, 2024

github-actions bot commented Dec 11, 2024

github-actions bot commented Dec 11, 2024

github-actions bot commented Dec 12, 2024

github-actions bot commented Dec 13, 2024

github-actions bot commented Dec 18, 2024

github-actions bot commented Dec 27, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

github-actions bot commented Dec 28, 2024

ndgrigorian commented Dec 28, 2024

oleksandr-pavlyk commented Dec 28, 2024

ndgrigorian commented Jan 5, 2025

github-actions bot commented Jan 5, 2025

github-actions bot commented Jan 5, 2025

Implements `top_k` in dpctl.tensor #1921

Implements `top_k` in dpctl.tensor #1921

ndgrigorian commented Dec 4, 2024 •

edited

Loading

github-actions bot commented Dec 4, 2024 •

edited

Loading

coveralls commented Dec 4, 2024 •

edited

Loading