Specify access mode in __result_and_scratch_storage methods
#1909
+55
−20
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 <sergey.kopienko@intel.com>
…add template parameter sycl::access_mode into __result_and_scratch_storage::__get_result_acc Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
…add template parameter sycl::access_mode into __result_and_scratch_storage::__get_scratch_acc Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
Signed-off-by: Sergey Kopienko <sergey.kopienko@intel.com>
SergeyKopienko
commented
Oct 17, 2024
| @@ -319,6 +319,8 @@ struct __parallel_scan_submitter<_CustomName, __internal::__optional_kernel_name | |||
| // 1. Local scan on each workgroup | |||
| auto __submit_event = __exec.queue().submit([&](sycl::handler& __cgh) { | |||
| oneapi::dpl::__ranges::__require_access(__cgh, __rng1, __rng2); //get an access to data under SYCL buffer | |||
| // TODO what is correct access mode here for __get_scratch_acc call? | |||
There was a problem hiding this comment.
@danhoeflinger do you know the correct template param state here and in other TODO's ?
There was a problem hiding this comment.
i think this can be write
| @@ -478,12 +478,11 @@ using __repacked_tuple_t = typename __repacked_tuple<T>::type; | |||
| template <typename _ContainerOrIterable> | |||
| using __value_t = typename __internal::__memobj_traits<_ContainerOrIterable>::value_type; | |||
|
|
|||
| template <typename _T> | |||
| template <typename _T, sycl::access_mode _AccessMode = sycl::access_mode::read_write> | |||
There was a problem hiding this comment.
I don't think we gain anything from this default template argument. All the usages specify explicitly at this level.
Comment on lines
+518
to
+531
| static auto __get_access_mode_tag() | ||
| { | ||
| if constexpr (_AccessMode == sycl::access::mode::read) | ||
| return sycl::read; | ||
|
|
||
| else if constexpr (_AccessMode == sycl::access::mode::write) | ||
| return sycl::write; | ||
|
|
||
| else if constexpr (_AccessMode == sycl::access::mode::read_write) | ||
| return sycl::read_write; | ||
|
|
||
| else | ||
| static_assert(false, "Unknown _AccessMode state"); | ||
| } |
There was a problem hiding this comment.
Should this be a general utility function outside of this structure (with a template arg)? It's useful in the structure but it seems to have general applicability. I'm a little surprised this utility doesn't exist in the specification itself actually.
Also the function itself should be able to be constexpr.
| @@ -453,8 +454,8 @@ struct __parallel_reduce_then_scan_scan_submitter< | |||
| __dpl_sycl::__local_accessor<_InitValueType> __sub_group_partials(__num_sub_groups_local + 1, __cgh); | |||
| __cgh.depends_on(__prior_event); | |||
| oneapi::dpl::__ranges::__require_access(__cgh, __in_rng, __out_rng); | |||
| auto __temp_acc = __scratch_container.__get_scratch_acc(__cgh); | |||
| auto __res_acc = __scratch_container.__get_result_acc(__cgh); | |||
| auto __temp_acc = __scratch_container.__get_scratch_acc<sycl::access_mode::read_write>(__cgh); | |||
There was a problem hiding this comment.
I think this can be just read access.
Comment on lines
+304
to
+305
| _InitValueType* __temp_ptr = | ||
| _TmpStorageAcc::__get_usm_or_buffer_accessor_ptr<sycl::access_mode::write>(__temp_acc); |
There was a problem hiding this comment.
I don't think we need an access mode here.
Comment on lines
+386
to
+387
| __result_and_scratch_storage_t::__get_usm_or_buffer_accessor_ptr<sycl::access_mode::write>( | ||
| __res_acc, 2 * __n_groups); |
There was a problem hiding this comment.
i dont think we need an access mode here.
Comment on lines
+256
to
+257
| // TODO what is correct access mode here for __get_scratch_acc call? | ||
| // Is default sycl::access_mode::read_write is ok? |
There was a problem hiding this comment.
I think this is read access.
| @@ -319,6 +319,8 @@ struct __parallel_scan_submitter<_CustomName, __internal::__optional_kernel_name | |||
| // 1. Local scan on each workgroup | |||
| auto __submit_event = __exec.queue().submit([&](sycl::handler& __cgh) { | |||
| oneapi::dpl::__ranges::__require_access(__cgh, __rng1, __rng2); //get an access to data under SYCL buffer | |||
| // TODO what is correct access mode here for __get_scratch_acc call? | |||
There was a problem hiding this comment.
i think this can be write
| @@ -340,6 +342,8 @@ struct __parallel_scan_submitter<_CustomName, __internal::__optional_kernel_name | |||
| auto __iters_per_single_wg = oneapi::dpl::__internal::__dpl_ceiling_div(__n_groups, __wgroup_size); | |||
| __submit_event = __exec.queue().submit([&](sycl::handler& __cgh) { | |||
| __cgh.depends_on(__submit_event); | |||
| // TODO what is correct access mode here for __get_scratch_acc call? | |||
| // Is default sycl::access_mode::read_write is ok? | |||
| @@ -362,7 +366,11 @@ struct __parallel_scan_submitter<_CustomName, __internal::__optional_kernel_name | |||
| auto __final_event = __exec.queue().submit([&](sycl::handler& __cgh) { | |||
| __cgh.depends_on(__submit_event); | |||
| oneapi::dpl::__ranges::__require_access(__cgh, __rng1, __rng2); //get an access to data under SYCL buffer | |||
| // TODO what is correct access mode here for __get_scratch_acc call? | |||
| // Is default sycl::access_mode::read_write is ok? | |||
Comment on lines
+372
to
+373
| // TODO what is right access mode here? | ||
| // sycl::access_mode::read_write is used by default. |
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.