Add iota kernel #1946

oleksandr-pavlyk · 2024-12-20T22:56:00Z

This PR builds on top of feature/topk branch.

It adds iota_impl in new sort_utils.hpp file, and uses it in merge_sort.hpp, radix_sort.hpp and topk.hpp.

It also fixes possible USM allocation leak in exception handling.

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?

…d_largest test

Reuse that function call in sorting code-base where argsort is used.

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp

github-actions · 2024-12-20T23:31:07Z

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

github-actions · 2024-12-20T23:39:15Z

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

github-actions · 2024-12-22T21:07:33Z

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

github-actions · 2024-12-22T21:38:56Z

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

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp

github-actions · 2024-12-23T05:15:21Z

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

github-actions · 2024-12-24T00:57:25Z

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

…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...] ```

github-actions · 2024-12-24T16:19:06Z

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

github-actions · 2024-12-24T16:35:41Z

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

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp

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

github-actions · 2024-12-25T13:24:40Z

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

github-actions · 2024-12-25T13:25:27Z

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

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

github-actions · 2024-12-25T17:49:36Z

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

ndgrigorian · 2024-12-25T21:30:20Z

Ping @AlexanderKalistratov
If the utilities added here (smart_malloc_* functions) are merged in, they can be reused by dpnp and replace the similar code in choose PR

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.

github-actions · 2024-12-26T00:10:48Z

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

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp

Add check of computed against expected indices

github-actions · 2024-12-27T03:47:21Z

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

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

github-actions · 2024-12-27T15:42:12Z

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

oleksandr-pavlyk · 2024-12-27T18:18:12Z

I suggest we exclude these failing dpt.int8, and dpt.int16 tests in the workflow, now that we know they are caused by a CPU driver issue for certain architectures.

ndgrigorian

This LGTM, we can merge this into the topk branch and drop the test file PR, then remove the commit that adds test_top_k_largest_1d_radix_i1

oleksandr-pavlyk and others added 4 commits December 13, 2024 18:48

Add test file for top_k functionality

be9f5ce

Adds test to triage AMD CPU top_k failure

ff81365

radix sort triage test now parametrizes for identical sizes to topk_1…

58da416

…d_largest test

Add sort_utils.hpp with iota_impl

fd4bab3

Reuse that function call in sorting code-base where argsort is used.

oleksandr-pavlyk requested a review from ndgrigorian as a code owner December 20, 2024 22:56

ndgrigorian reviewed Dec 20, 2024

View reviewed changes

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp Outdated Show resolved Hide resolved

oleksandr-pavlyk force-pushed the add-iota-kernel branch from 29d7198 to f1b2045 Compare December 23, 2024 00:38

ndgrigorian reviewed Dec 23, 2024

View reviewed changes

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp Outdated Show resolved Hide resolved

ndgrigorian reviewed Dec 23, 2024

View reviewed changes

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp Outdated Show resolved Hide resolved

ndgrigorian reviewed Dec 23, 2024

View reviewed changes

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp Outdated Show resolved Hide resolved

oleksandr-pavlyk force-pushed the add-iota-kernel branch from f1b2045 to 51ead2b Compare December 23, 2024 04:30

ndgrigorian reviewed Dec 23, 2024

View reviewed changes

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp Show resolved Hide resolved

oleksandr-pavlyk added 3 commits December 24, 2024 09:48

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

d9cc919

…gid)

Add explicit read into registers

eb28d1a

oleksandr-pavlyk force-pushed the add-iota-kernel branch from bbb55f1 to da3fbcc Compare December 24, 2024 15:52

ndgrigorian reviewed Dec 25, 2024

View reviewed changes

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp Outdated Show resolved Hide resolved

oleksandr-pavlyk added 2 commits December 25, 2024 06:41

Factor out write-out kernel into separate detail function

79feecd

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

Typo: smart_malloc_jost -> smart_malloc_host

83b1771

Reimplemented map_back_impl to process few elements per work-item

1c25d0b

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)

a74c788

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.

ndgrigorian reviewed Dec 26, 2024

View reviewed changes

dpctl/tensor/libtensor/include/kernels/sorting/topk.hpp Outdated Show resolved Hide resolved

ndgrigorian reviewed Dec 26, 2024

View reviewed changes

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp Show resolved Hide resolved

ndgrigorian reviewed Dec 27, 2024

View reviewed changes

dpctl/tensor/libtensor/include/utils/sycl_alloc_utils.hpp Show resolved Hide resolved

oleksandr-pavlyk added 2 commits December 26, 2024 21:03

Remove dead code disable via preprocessor conditional

9294095

Add information displayed on failure, renamed variables

1bb83bc

Add check of computed against expected indices

oleksandr-pavlyk added 2 commits December 27, 2024 08:56

Add comment to explain ceil_log2 algo

fde7408

Add static assers to async_smart_free

4c8d7d7

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

ndgrigorian approved these changes Dec 27, 2024

View reviewed changes

oleksandr-pavlyk merged commit e161ea1 into feature/topk Dec 27, 2024
41 of 52 checks passed

oleksandr-pavlyk deleted the add-iota-kernel branch December 27, 2024 18:57

oleksandr-pavlyk mentioned this pull request Dec 27, 2024

Triage top_k test failure from radix-sort-based implementation #1945

Closed

8 tasks

Add iota kernel #1946

Add iota kernel #1946

Uh oh!

Conversation

oleksandr-pavlyk commented Dec 20, 2024

Uh oh!

Uh oh!

github-actions bot commented Dec 20, 2024 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions bot commented Dec 20, 2024

Uh oh!

github-actions bot commented Dec 22, 2024

Uh oh!

github-actions bot commented Dec 22, 2024

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

github-actions bot commented Dec 23, 2024

Uh oh!

github-actions bot commented Dec 24, 2024

Uh oh!

github-actions bot commented Dec 24, 2024

Uh oh!

github-actions bot commented Dec 24, 2024

Uh oh!

Uh oh!

github-actions bot commented Dec 25, 2024

Uh oh!

github-actions bot commented Dec 25, 2024

Uh oh!

github-actions bot commented Dec 25, 2024

Uh oh!

ndgrigorian commented Dec 25, 2024

Uh oh!

github-actions bot commented Dec 26, 2024

Uh oh!

Uh oh!

Uh oh!

Uh oh!

github-actions bot commented Dec 27, 2024

Uh oh!

github-actions bot commented Dec 27, 2024

Uh oh!

oleksandr-pavlyk commented Dec 27, 2024

Uh oh!

ndgrigorian left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

github-actions bot commented Dec 20, 2024 •

edited

Loading