-
Notifications
You must be signed in to change notification settings - Fork 114
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
Do not use deprecated sub-group load/store extension #1979
Conversation
fee39ab
to
0c0eae5
Compare
@@ -598,7 +602,7 @@ struct __parallel_copy_if_static_single_group_submitter<_Size, _ElemsPerItem, _W | |||
const ::std::uint16_t __subgroup_id = __subgroup.get_group_id(); | |||
const ::std::uint16_t __subgroup_size = __subgroup.get_local_linear_range(); | |||
|
|||
#if _ONEDPL_SYCL_SUB_GROUP_LOAD_STORE_PRESENT | |||
#if _ONEDPL_LIBSYCL_SUB_GROUP_LOAD_STORE_PRESENT |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
-
This is a "logic mix" of preprocessor code and C++ code.... It looks a mess..
I would recommend "take a look/imagine" on C++ code which remains after the preprocessor keeps one branch "#if" and the other branch "#else" and make the code clean-up (here and everywhere in that PR). -
BTW,
_ONEDPL_LIBSYCL_SUB_GROUP_LOAD_STORE_PRESENT
is defined as identical zero.
https://github.com/oneapi-src/oneDPL/pull/1979/files?diff=split&w=0#diff-521f7db06a55567c1a1dffa855dde585fb6bc5fbe8633d19b528356e3a501ea0R91
M.b. it makes sense to remove the branches_ONEDPL_LIBSYCL_SUB_GROUP_LOAD_STORE_PRESENT
at all?
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I've removed all the branches. Initially, they were there to preserve the logic if we decide to add newer versions of sub-group load/store. If we settle with this, I can create an issue to remember this optimization opportunity.
Removing the sub-group load store logic seems fine to me as it is currently unused. It may be worth creating an issue to investigate the performance difference and determine if it is worth it to adopt. I suggest we wait for @adamfidel's input when he returns since I believe he introduced sub-group load / store path here. |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Thanks for cleaning this up. The idea was to come back and re-enable this code path at a later point, but since this has been dead code for some time now, I agree it makes sense to just remove it for now.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
Code sections which use
sycl::sub_group::<load()|store()>
DPC++ extension are not entirely disabled. Non-template expressions are still checked by a compiler even if they are within a discardedif constexpr
branch.This makes the code non-compatible with compilers other than DPC++, and it results in a deprecation warning with the most recent DPC++ builds.
A mock-up example: https://godbolt.org/z/svhnY14r5