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.

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

[KT] Remove data() function from range view classes #1044

Conversation

SergeyKopienko
Copy link
Contributor

@SergeyKopienko SergeyKopienko commented Aug 9, 2023

This PR created for discussion how to avoid two new data() methods in esimd radix sort code.

@SergeyKopienko
Copy link
Contributor Author

SergeyKopienko commented Aug 9, 2023

for additional details please see: https://www.intel.com/content/www/us/en/docs/dpcpp-cpp-compiler/developer-guide-reference/2023-0/explicit-simd-sycl-extension.html

Unsupported standard SYCL APIs:
- Local accessors. Local memory is allocated and accessed via explicit device-side APIs.
- 2D and 3D accessors.
- Constant accessors.
- sycl::accessor::get_pointer(). All memory accesses through an accessor are done via explicit APIs. Example: sycl::ext::intel::esimd::block_store(acc, offset)
- Accessors with offsets and/or access range specified.
sycl::sampler and sycl::stream classes.

@SergeyKopienko
Copy link
Contributor Author

@MikeDvorskiy what do you think about this?

@@ -245,6 +245,9 @@ class guard_view
// Unified funciton to get pointer or accessor to use inside ESIMD kernels
_Iterator data()
Copy link
Contributor Author

@SergeyKopienko SergeyKopienko Aug 9, 2023

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@MikeDvorskiy, I believe we are able to remove this function only.

@@ -100,6 +100,8 @@ class all_view
// Unified funciton to get pointer or accessor to use inside ESIMD kernels
__accessor_t data()
Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@SergeyKopienko SergeyKopienko marked this pull request as ready for review August 9, 2023 15:48
@MikeDvorskiy
Copy link
Contributor

@MikeDvorskiy what do you think about this?

Why begin method is not enough?

@SergeyKopienko
Copy link
Contributor Author

@MikeDvorskiy, the begin() isn't enough due existing restrictions from https://www.intel.com/content/www/us/en/docs/dpcpp-cpp-compiler/developer-guide-reference/2023-0/explicit-simd-sycl-extension.html : we are unable to work with data address from sycl::accessor together with esimd implementation. @dmitriy-sobolev early tried to do this.

@SergeyKopienko SergeyKopienko force-pushed the dev/skopienko/esimd/esimd-radix-sort_onesweep_extend_ranges branch from 50d2d68 to e920601 Compare August 9, 2023 16:51
@dmitriy-sobolev dmitriy-sobolev changed the title Remove data() function from range view classes [KT] Remove data() function from range view classes Aug 9, 2023
@SergeyKopienko SergeyKopienko force-pushed the dev/skopienko/esimd/esimd-radix-sort_onesweep_extend_ranges branch from e920601 to fe2cc72 Compare August 14, 2023 12:48
@SergeyKopienko
Copy link
Contributor Author

@MikeDvorskiy finally, are you agree with current implementation on data() functions in esimd radix sort code?

@SergeyKopienko SergeyKopienko force-pushed the dev/skopienko/esimd/esimd-radix-sort_onesweep_extend_ranges branch 9 times, most recently from 7c5b2cd to 9abd10b Compare August 21, 2023 14:26
@SergeyKopienko SergeyKopienko force-pushed the dev/skopienko/esimd/esimd-radix-sort_onesweep_extend_ranges branch from 9abd10b to 8b9ed3f Compare August 29, 2023 07:27
@SergeyKopienko SergeyKopienko force-pushed the dev/skopienko/esimd/esimd-radix-sort_onesweep_extend_ranges branch from 8b9ed3f to 3dc42ec Compare August 30, 2023 08:19
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 2 times, most recently from 8500608 to 56e1870 Compare September 1, 2023 12:31
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 7 times, most recently from af77e4d to f48fefe Compare September 19, 2023 07:45
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 2 times, most recently from ffc3dfb to 8726ac4 Compare September 22, 2023 12:41
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 3 times, most recently from f77a660 to 5aba75e Compare October 5, 2023 16:00
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 4 times, most recently from 3161b3e to 65a6b34 Compare October 12, 2023 08:58
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 3 times, most recently from 6a86bbe to 8a3bf7b Compare October 19, 2023 18:34
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 3 times, most recently from c0bce0b to a5826c5 Compare October 25, 2023 08:36
@dmitriy-sobolev dmitriy-sobolev force-pushed the esimd-radix-sort branch 2 times, most recently from a9ca5d4 to 9eb5e27 Compare November 1, 2023 09:13
@SergeyKopienko SergeyKopienko force-pushed the esimd-radix-sort branch 2 times, most recently from b85452d to 65d10e9 Compare November 13, 2023 08:31
@SergeyKopienko SergeyKopienko deleted the dev/skopienko/esimd/esimd-radix-sort_onesweep_extend_ranges branch November 13, 2023 10:47
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants