Improvement to performance of tensor.sum #1303

oleksandr-pavlyk · 2023-07-25T15:34:09Z

Transition sum-reduction from sycl::nd_range<2> to sycl::nd_range<1>

This improves performance 8x-fold:

In [1]: import dpctl.tensor as dpt

In [2]: x = dpt.ones((4096, 4096), dtype="f4")

In [3]: y = dpt.sum(x, axis=0)

In [4]: %time y = dpt.sum(x, axis=0)
CPU times: user 2.64 ms, sys: 4.4 ms, total: 7.04 ms
Wall time: 10 ms

In [5]: %time y = dpt.sum(x, axis=0)
CPU times: user 1.93 ms, sys: 3.22 ms, total: 5.16 ms
Wall time: 4.74 ms

In [6]: %time y = dpt.sum(x, axis=0)
CPU times: user 1.7 ms, sys: 2.83 ms, total: 4.53 ms
Wall time: 4.1 ms

In [7]: %time y = dpt.sum(x, axis=0)
CPU times: user 1.98 ms, sys: 3.3 ms, total: 5.28 ms
Wall time: 4.7 ms

The timing before was around 38ms

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?
If this PR is a work in progress, are you opening the PR as a draft?

github-actions · 2023-07-25T16:04:57Z

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

coveralls · 2023-07-25T16:06:27Z

coverage: 84.95%. remained the same when pulling 9f54428 on reduction-changes into 074ec3a on master.

github-actions · 2023-07-25T16:34:32Z

Array API standard conformance tests for dpctl=0.14.6dev0=py310h7bf5fec_28 ran successfully.
Passed: 477
Failed: 523
Skipped: 119

ndgrigorian · 2023-07-26T03:48:51Z

Currently giving incorrect results for some cases with axes:

In [1]: import dpctl.tensor as dpt, numpy as np

In [2]: x = dpt.ones((2, 4098, 4098), dtype="i2")

In [3]: x_np = dpt.asnumpy(x)

In [4]: dpt.sum(x, axis=(0, 1))
Out[4]: usm_ndarray([8704, 8704, 8704, ...,  68,   68,   68])

In [5]: np.sum(x_np, axis=(0, 1))
Out[5]: array([8196, 8196, 8196, ..., 8196, 8196, 8196])

This improves performance 8x-fold: ``` In [1]: import dpctl.tensor as dpt In [2]: x = dpt.ones((4096, 4096), dtype="f4") In [3]: y = dpt.sum(x, axis=0) In [4]: %time y = dpt.sum(x, axis=0) CPU times: user 2.64 ms, sys: 4.4 ms, total: 7.04 ms Wall time: 10 ms In [5]: %time y = dpt.sum(x, axis=0) CPU times: user 1.93 ms, sys: 3.22 ms, total: 5.16 ms Wall time: 4.74 ms In [6]: %time y = dpt.sum(x, axis=0) CPU times: user 1.7 ms, sys: 2.83 ms, total: 4.53 ms Wall time: 4.1 ms In [7]: %time y = dpt.sum(x, axis=0) CPU times: user 1.98 ms, sys: 3.3 ms, total: 5.28 ms Wall time: 4.7 ms ``` The timing before was around 38ms

github-actions · 2023-07-28T20:15:07Z

Array API standard conformance tests for dpctl=0.14.6dev0=py310h7bf5fec_53 ran successfully.
Passed: 567
Failed: 433
Skipped: 119

- Adjusted to reduce branching and hopefully improve vectorization of the loop by removing a conditional

github-actions · 2023-07-30T06:09:05Z

Array API standard conformance tests for dpctl=0.14.6dev0=py310h7bf5fec_62 ran successfully.
Passed: 594
Failed: 406
Skipped: 119

github-actions · 2023-08-11T17:55:08Z

Array API standard conformance tests for dpctl=0.14.6dev2=py310h7bf5fec_7 ran successfully.
Passed: 913
Failed: 87
Skipped: 119

github-actions · 2023-08-13T14:42:12Z

Array API standard conformance tests for dpctl=0.14.6dev2=py310h7bf5fec_10 ran successfully.
Passed: 913
Failed: 87
Skipped: 119

oleksandr-pavlyk · 2023-08-18T18:24:09Z

The timing now:

(dev_dpctl) opavlyk@opavlyk-mobl:~/repos/dpctl$ ipython
Python 3.9.12 (main, Jun  1 2022, 11:38:51)
Type 'copyright', 'credits' or 'license' for more information
IPython 8.14.0 -- An enhanced Interactive Python. Type '?' for help.

In [1]: import dpctl.tensor as dpt

In [2]: x = dpt.ones((4096, 4096), dtype="f4")

In [3]: y = dpt.sum(x, axis=0)

In [4]: %time y = dpt.sum(x, axis=0)
CPU times: user 25.2 ms, sys: 6.65 ms, total: 31.9 ms
Wall time: 33.8 ms

In [5]: %time y = dpt.sum(x, axis=0)
CPU times: user 10.6 ms, sys: 20.1 ms, total: 30.7 ms
Wall time: 30 ms

In [6]: %time y = dpt.sum(x, axis=0)
CPU times: user 13.4 ms, sys: 19 ms, total: 32.4 ms
Wall time: 31.1 ms

In [7]: %time y = dpt.sum(x, axis=0)
CPU times: user 10.4 ms, sys: 20.1 ms, total: 30.4 ms
Wall time: 29.5 ms

In [8]: !git describe --tag
0.14.6dev2-7-g9f544288f

It still has improved over 38 ms, but my claim of reducing to 4.7ms was erroneous (it was faster, due to it being incorrect).

@ndgrigorian I would prefer to merge this change, and defer further work to later PRs.

ndgrigorian · 2023-08-18T18:27:41Z

The timing now:

(dev_dpctl) opavlyk@opavlyk-mobl:~/repos/dpctl$ ipython
Python 3.9.12 (main, Jun  1 2022, 11:38:51)
Type 'copyright', 'credits' or 'license' for more information
IPython 8.14.0 -- An enhanced Interactive Python. Type '?' for help.

In [1]: import dpctl.tensor as dpt

In [2]: x = dpt.ones((4096, 4096), dtype="f4")

In [3]: y = dpt.sum(x, axis=0)

In [4]: %time y = dpt.sum(x, axis=0)
CPU times: user 25.2 ms, sys: 6.65 ms, total: 31.9 ms
Wall time: 33.8 ms

In [5]: %time y = dpt.sum(x, axis=0)
CPU times: user 10.6 ms, sys: 20.1 ms, total: 30.7 ms
Wall time: 30 ms

In [6]: %time y = dpt.sum(x, axis=0)
CPU times: user 13.4 ms, sys: 19 ms, total: 32.4 ms
Wall time: 31.1 ms

In [7]: %time y = dpt.sum(x, axis=0)
CPU times: user 10.4 ms, sys: 20.1 ms, total: 30.4 ms
Wall time: 29.5 ms

In [8]: !git describe --tag
0.14.6dev2-7-g9f544288f

It still has improved over 38 ms, but my claim of reducing to 4.7ms was erroneous (it was faster, due to it being incorrect).

@ndgrigorian I would prefer to merge this change, and defer further work to later PRs.

I agree. I'll approve this PR.

github-actions · 2023-08-18T18:29:25Z

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

github-actions · 2023-08-18T19:40:55Z

Array API standard conformance tests for dpctl=0.14.6dev3=py310ha25a700_42 ran successfully.
Passed: 913
Failed: 87
Skipped: 119

oleksandr-pavlyk requested review from antonwolfy, ndgrigorian and npolina4 July 25, 2023 15:34

oleksandr-pavlyk added 3 commits July 28, 2023 06:14

Replaced nd_range<2> with nd_range<1> for non-atomic case as well

2f79acb

Add test based on example from @ndgrigorian's feedback to PR

5c4f980

oleksandr-pavlyk force-pushed the reduction-changes branch from 1d5228a to 5c4f980 Compare July 28, 2023 19:10

ndgrigorian added 2 commits July 29, 2023 21:47

Boolean reductions transitioned from nd_range<2> to nd_range<1>

a5aee5b

Strided boolean reduction loop tweak

da1cc5b

- Adjusted to reduce branching and hopefully improve vectorization of the loop by removing a conditional

Merge remote-tracking branch 'origin/master' into reduction-changes

63b2799

Specify name for the atomic reduction initialization kernel

9f54428

ndgrigorian approved these changes Aug 18, 2023

View reviewed changes

oleksandr-pavlyk merged commit b008b8b into master Aug 18, 2023

oleksandr-pavlyk deleted the reduction-changes branch August 18, 2023 18:29

oleksandr-pavlyk mentioned this pull request Aug 18, 2023

Development milestone 0.14.6dev4 #1354

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

Improvement to performance of tensor.sum #1303

Improvement to performance of tensor.sum #1303

Uh oh!

oleksandr-pavlyk commented Jul 25, 2023

Uh oh!

github-actions bot commented Jul 25, 2023

Uh oh!

coveralls commented Jul 25, 2023 •

edited

Loading

Uh oh!

github-actions bot commented Jul 25, 2023

Uh oh!

ndgrigorian commented Jul 26, 2023 •

edited

Loading

Uh oh!

github-actions bot commented Jul 28, 2023

Uh oh!

github-actions bot commented Jul 30, 2023

Uh oh!

github-actions bot commented Aug 11, 2023

Uh oh!

github-actions bot commented Aug 13, 2023

Uh oh!

oleksandr-pavlyk commented Aug 18, 2023

Uh oh!

ndgrigorian commented Aug 18, 2023

Uh oh!

github-actions bot commented Aug 18, 2023

Uh oh!

github-actions bot commented Aug 18, 2023

Uh oh!

Uh oh!

Improvement to performance of tensor.sum #1303

Improvement to performance of tensor.sum #1303

Uh oh!

Conversation

oleksandr-pavlyk commented Jul 25, 2023

Uh oh!

github-actions bot commented Jul 25, 2023

Uh oh!

coveralls commented Jul 25, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions bot commented Jul 25, 2023

Uh oh!

ndgrigorian commented Jul 26, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

github-actions bot commented Jul 28, 2023

Uh oh!

github-actions bot commented Jul 30, 2023

Uh oh!

github-actions bot commented Aug 11, 2023

Uh oh!

github-actions bot commented Aug 13, 2023

Uh oh!

oleksandr-pavlyk commented Aug 18, 2023

Uh oh!

ndgrigorian commented Aug 18, 2023

Uh oh!

github-actions bot commented Aug 18, 2023

Uh oh!

github-actions bot commented Aug 18, 2023

Uh oh!

Uh oh!

coveralls commented Jul 25, 2023 •

edited

Loading

ndgrigorian commented Jul 26, 2023 •

edited

Loading