Skip to content
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

Jump to bottom

[DO NOT MERGE] Sasha triage topk test failure amd #1947

Closed
wants to merge 25 commits into from
Closed

[DO NOT MERGE] Sasha triage topk test failure amd #1947

wants to merge 25 commits into from

Conversation

oleksandr-pavlyk
Copy link
Collaborator

  • 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?

ndgrigorian and others added 13 commits December 11, 2024 10:46
@ndgrigorian
Implements top_k
d03579f 
The implementation leverages existing merge-sort code, and partially sorts the array in cases where a parial sort reduces the size of temporary memory allocation
@ndgrigorian
Add top_k to doc sources
3955f49
@ndgrigorian
Use std::size_t for k instead of py::ssize_t
2e0859c 
Reduces amount of casting. `k` will need to fit in `py::ssize_t` regardless.
@ndgrigorian
Add implementation of top_k using radix sort
26718f3
@ndgrigorian
Reduce code duplication by using std::optional in py_topk directly
6c43a4b 
Instead of using an overload to handle the `axis=None` case, use std::optional and check for trailing_dims_to_search in validation logic
@ndgrigorian
Clean up top_k docstring
442f96f
@ndgrigorian
Move top_k into _sorting.py
3f15cf8
@ndgrigorian
top_k raises when k > 1 and input is a scalar
5b0b80f
@ndgrigorian
Fix bug in top_k partial merge sort implementation
8f38b80 
rounded value of k must be divisible by the merge sort chunk size
@oleksandr-pavlyk
Add test file for top_k functionality
be9f5ce
@ndgrigorian
Adds test to triage AMD CPU top_k failure
ff81365
@ndgrigorian
radix sort triage test now parametrizes for identical sizes to topk_1…
58da416 
…d_largest test
@oleksandr-pavlyk
Add sort_utils.hpp with iota_impl
fd4bab3 
Reuse that function call in sorting code-base where argsort is used.
@oleksandr-pavlyk oleksandr-pavlyk changed the title Sasha triage topk test failure amd [DO NOT MERGE] Sasha triage topk test failure amd Dec 22, 2024
@github-actions GitHub Actions
Copy link

View rendered docs @ https://intelpython.github.io/dpctl/pulls/1947/index.html

@github-actions GitHub Actions
Copy link

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

@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch 3 times, most recently from de53055 to 3e5e303 Compare December 23, 2024 13:24
@github-actions GitHub Actions
Copy link

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

@coveralls
Copy link
Collaborator

Coverage Status

coverage: 87.488% (-0.2%) from 87.659%
when pulling 3e5e303 on sasha-triage-topk-test-failure-amd
into 678b4cf on master.

@oleksandr-pavlyk oleksandr-pavlyk changed the base branch from master to add-iota-kernel December 23, 2024 17:34
@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from 3e5e303 to d50092a Compare December 23, 2024 17:36
@github-actions GitHub Actions
Copy link

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

@oleksandr-pavlyk
Simplify write-out kernels in topk implementation (avoid recomputing …
d9cc919 
…gid)
@oleksandr-pavlyk
Add explicit read into registers
eb28d1a
@oleksandr-pavlyk
Use unique_ptr as temporary owner of USM allocation
da3fbcc 
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 GitHub Actions
Copy link

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

@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from d50092a to b411407 Compare December 24, 2024 17:17
@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from 8214855 to 04b6629 Compare December 25, 2024 12:44
@github-actions GitHub Actions
Copy link

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

@github-actions GitHub Actions
Copy link

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

@oleksandr-pavlyk
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.
@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from 04b6629 to dfb521f Compare December 25, 2024 17:09
@github-actions GitHub Actions
Copy link

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

@github-actions GitHub Actions
Copy link

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

@github-actions GitHub Actions
Copy link

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

@oleksandr-pavlyk
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.
@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from c351d0b to c1f8a74 Compare December 25, 2024 23:29
@github-actions GitHub Actions
Copy link

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

@github-actions GitHub Actions
Copy link

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

@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from 0869128 to 210500f Compare December 26, 2024 14:26
@github-actions GitHub Actions
Copy link

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

@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from 210500f to 387a3d9 Compare December 27, 2024 03:08
@github-actions GitHub Actions
Copy link

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

@oleksandr-pavlyk
Add comment to explain ceil_log2 algo
fde7408
@oleksandr-pavlyk
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.
@oleksandr-pavlyk
Run test_usm_ndarray_topk under gdb in few combination
fd65511
@oleksandr-pavlyk oleksandr-pavlyk force-pushed the sasha-triage-topk-test-failure-amd branch from 387a3d9 to fd65511 Compare December 27, 2024 14:58
@github-actions GitHub Actions
Copy link

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

Base automatically changed from add-iota-kernel to feature/topk December 27, 2024 18:57
@ndgrigorian ndgrigorian force-pushed the feature/topk branch 4 times, most recently from 84d1388 to 809cb70 Compare December 28, 2024 01:13
@oleksandr-pavlyk oleksandr-pavlyk deleted the sasha-triage-topk-test-failure-amd branch January 8, 2025 22:48
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@ndgrigorian ndgrigorian Awaiting requested review from ndgrigorian

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
3 participants
@oleksandr-pavlyk @coveralls @ndgrigorian