[SYCL][SubBlockNDRange] Extend spec to allow dealing with USM pointers

[bjoo]

Hi, following conversations with @jbrodman and @Pennycook I am leaving this issue here as a placeholder. Currently the SubGroupNDRange Load and Save functions use sycl::multi_ptr<>
It would be good if they could accept raw USM pointers, both to connect up the USM and SubGroupNDRange extensions and also to provide for the possibility of casting. The use case for casting is in the case of SIMD operations. It may make sense allocate buffers based on SIMD Classes e.g. sycl::vec<T,N>. In the case of SYCL buffers this means that the accessors to the buffer are of type sycl::accessor<vec<T,N>,...> and the pointers returned by the accessors get_pointer() method (which convert to multi_ptr<>) are vec<T,N>* However, for the load and save of SIMD data
a regular multi_ptr or T* (USM) would be more appropriate. Accessors cannot be cast, and neither can multi_ptr<> as far as I know. However USM pointers can. In that case if one allocated with a USM appropriate malloc( N * vec<T,N> ) and indexed the array in units of vec<T,N> it would always be easy to cast the pointer to type T*

[Pennycook]

I've tested this since #1183 and it almost works. You can now write something like sg.store(global_ptr<float>(usm_ptr), x), but you can't write something like sg.store(usm_ptr, x). The latter still fails despite pointers now being implicitly convertible to multi_ptr, because the type and space of the multi_ptr template arguments need to be deduced.

I'll need to think about this some more; please let me know if anybody has any ideas about how to solve this issue in a straightforward way.

[AlexeySachkov]

I think that this functionality should be available since #6893

However, we have since then deprecated our own sub-group extension in favor of the core SYCL 2020 functionality and load/store methods are in a weird state: they are not documented by the core SYCL specification and we don't have a sub_group class extension for them anymore.

I think this will be resolved by sycl_ext_oneapi_group_load_store extension proposed in #7593.

[github-actions]

Hi! There have been no updates for at least the last 60 days, though the issue has assignee(s).

@jbrodman, could you please take one of the following actions:

• provide an update if you have any
• unassign yourself if you're not looking / going to look into this issue
• mark this issue with the 'confirmed' label if you have confirmed the problem/request and our team should work on it
• close the issue if it has been resolved
• take any other suitable action.

Thanks!

[github-actions]

Hi! There have been no updates for at least the last 60 days, though the issue has assignee(s).

@jbrodman, could you please take one of the following actions:

• provide an update if you have any
• unassign yourself if you're not looking / going to look into this issue
• mark this issue with the 'confirmed' label if you have confirmed the problem/request and our team should work on it
• close the issue if it has been resolved
• take any other suitable action.

Thanks!

[AlexeySachkov]

I see that group load/store functionality was implemented in #13043, so I consider this issue to be complete and I will close it.

If there are questions, or if you disagree with my judgment here, please feel free to comment and/or re-open the issue.