Specify access mode in __result_and_scratch_storage methods
#1909
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
…add template parameter sycl::access_mode into __usm_or_buffer_accessor Signed-off-by: Sergey Kopienko <[email protected]>
…add template parameter sycl::access_mode into __result_and_scratch_storage::__get_result_acc Signed-off-by: Sergey Kopienko <[email protected]>
…add template parameter sycl::access_mode into __result_and_scratch_storage::__get_scratch_acc Signed-off-by: Sergey Kopienko <[email protected]>
Signed-off-by: Sergey Kopienko <[email protected]>
Signed-off-by: Sergey Kopienko <[email protected]>
SergeyKopienko
commented
Oct 17, 2024
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
Outdated
Show resolved
Hide resolved
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
Outdated
Show resolved
Hide resolved
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h
Outdated
Show resolved
Hide resolved
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h
Outdated
Show resolved
Hide resolved
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce.h
Outdated
Show resolved
Hide resolved
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce.h
Outdated
Show resolved
Hide resolved
…cess mode for __get_result_acc in the __parallel_scan_submitter::operator() - __final_event Signed-off-by: Sergey Kopienko <[email protected]>
…cess mode for __get_scratch_acc in the __parallel_scan_submitter::operator() - __submit_event Signed-off-by: Sergey Kopienko <[email protected]>
…cess mode for __get_scratch_acc in the __parallel_scan_submitter::operator() - __submit_event Signed-off-by: Sergey Kopienko <[email protected]>
… fix review comment: fix extra access mode specification in __parallel_transform_reduce_impl::submit Signed-off-by: Sergey Kopienko <[email protected]>
… fix access mode for __get_result_acc in the __parallel_transform_reduce_impl::submit - __reduce_event Signed-off-by: Sergey Kopienko <[email protected]>
… fix access mode for __get_scratch_acc in the __parallel_transform_reduce_impl::submit - __reduce_event Signed-off-by: Sergey Kopienko <[email protected]>
… fix access mode for __get_scratch_acc in the __parallel_transform_reduce_work_group_kernel_submitter::operator() Signed-off-by: Sergey Kopienko <[email protected]>
…remove default value for template parameter _AccessMode at struct __usm_or_buffer_accessor Signed-off-by: Sergey Kopienko <[email protected]>
…n_scan.h - fix review comment: fix extra access mode specification in __parallel_reduce_then_scan_reduce_submitter::operator() Signed-off-by: Sergey Kopienko <[email protected]>
…n_scan.h - fix access mode for __get_scratch_acc in the __parallel_reduce_then_scan_scan_submitter::operator() Signed-off-by: Sergey Kopienko <[email protected]>
…declare __usm_or_buffer_accessor::__get_access_mode_tag() as constexpr Signed-off-by: Sergey Kopienko <[email protected]>
…fix error in the __get_result_acc implementation Signed-off-by: Sergey Kopienko <[email protected]>
…fix error in the __get_scratch_acc implementation Signed-off-by: Sergey Kopienko <[email protected]>
…_or_buffer_accessor
… fix access mode for __get_scratch_acc in the __parallel_transform_reduce_work_group_kernel_submitter::operator() Signed-off-by: Sergey Kopienko <[email protected]>
Signed-off-by: Sergey Kopienko <[email protected]>
…fix error in the __usm_or_buffer_accessor::__get_access_mode_tag() implementation Signed-off-by: Sergey Kopienko <[email protected]>
…acc calls Signed-off-by: Sergey Kopienko <[email protected]>
…_acc calls Signed-off-by: Sergey Kopienko <[email protected]>
…ed at all due we shouldn't construct sycl::asseccor twice in the __usm_or_buffer_accessor::__usm_or_buffer_accessor Signed-off-by: Sergey Kopienko <[email protected]>
Signed-off-by: Sergey Kopienko <[email protected]>
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
Outdated
Show resolved
Hide resolved
I think in reality most of the target cases we are targeting for performance have access to USM and this will therefore be no change to performance. We could test it on some devices which don't support USM or force it into non-USM mode for the purposes of benchmarking but I'm unsure how useful that is. |
|
I do think that this can indirectly help with performance in our target cases by getting us closer to #1906. Switching the existing temporary storage that we have in buffers over to usm where it is available generally does provide a performance improvement. This PR improves the infrastructure to better enable using it elsewhere. |
…fix error: the usages of __dpl_sycl::__no_init{} not quite correct
Signed-off-by: Sergey Kopienko <[email protected]>
…y __dpl_sycl::__no_init{} for sycl::access_mode::write
Signed-off-by: Sergey Kopienko <[email protected]>
… specify __dpl_sycl::__no_init{} for sycl::access_mode::write
Signed-off-by: Sergey Kopienko <[email protected]>
…n_scan.h - specify __dpl_sycl::__no_init{} for sycl::access_mode::write
Signed-off-by: Sergey Kopienko <[email protected]>
Signed-off-by: Sergey Kopienko <[email protected]>
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_reduce_then_scan.h
Outdated
Show resolved
Hide resolved
…n_scan.h - fix review comment: I think this shouldn't be no_init, we read before write. Signed-off-by: Sergey Kopienko <[email protected]>
mmichel11
reviewed
Oct 21, 2024
There was a problem hiding this comment.
I noticed a couple of compilation errors in the buffer fallback path. I think it would be good to manually test the following scenarios by temporarily modifying the macros / function returns:
- The most common case:
_ONEDPL_SYCL_UNIFIED_USM_BUFFER_PRESENT=1,__use_USM_host_allocations=true - USM device allocations but no host memory:
_ONEDPL_SYCL_UNIFIED_USM_BUFFER_PRESENT=1,__use_USM_host_allocations=false. Testing on a CUDA / HIP device would also take this path. - The buffer fallback case:
_ONEDPL_SYCL_UNIFIED_USM_BUFFER_PRESENT=0. This is where we would probably find all issues related to access modes.
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
Outdated
Show resolved
Hide resolved
include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h
Outdated
Show resolved
Hide resolved
…fix compile error in the __result_and_scratch_storage for case when _ONEDPL_SYCL_UNIFIED_USM_BUFFER_PRESENT is undefined Signed-off-by: Sergey Kopienko <[email protected]>
…fix compile error in the __result_and_scratch_storage for case when _ONEDPL_SYCL_UNIFIED_USM_BUFFER_PRESENT is undefined Signed-off-by: Sergey Kopienko <[email protected]>
Signed-off-by: Sergey Kopienko <[email protected]>
SergeyKopienko
force-pushed
the
dev/skopienko/introduce_access_mode_into_usm_or_buffer_accessor
branch
from
ac67c14 to
b634f44
Compare
…parametrize struct __usm_or_buffer_accessor only by accessor Signed-off-by: Sergey Kopienko <[email protected]>
|
@danhoeflinger, @mmichel11 Could you please take a look again? |
SergeyKopienko
deleted the
dev/skopienko/introduce_access_mode_into_usm_or_buffer_accessor
branch
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
In this PR we specify access mode in the two
__result_and_scratch_storagemethods:__result_and_scratch_storage::__get_result_acc__result_and_scratch_storage::__get_scratch_accThe main idea - to write
<sycl::access_mode::write>is more correct then<sycl::access_mode::read_write>when we only write some data and to write<sycl::access_mode::read>is more correct then<sycl::access_mode::read_write>when we only read some data.