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

Fix sub-group load/store extension disabling logic #1979

Open
wants to merge 1 commit into
base: main
Choose a base branch
from

Conversation

dmitriy-sobolev
Copy link
Contributor

@dmitriy-sobolev dmitriy-sobolev commented Dec 22, 2024

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 discarded if 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


Even though the extension is deprecated, it still may make sense to keep the surrounding logic in case it can be reused with #1976 vector load/store or with SYCL_EXT_ONEAPI_GROUP_LOAD_STORE extension.

Base automatically changed from dev/dmitriy-sobolev/sycl2020-fallback to main December 23, 2024 10:55
@dmitriy-sobolev dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/sg-load branch from fee39ab to 0c0eae5 Compare December 23, 2024 11:53
@dmitriy-sobolev dmitriy-sobolev added this to the 2022.8.0 milestone Dec 24, 2024
@@ -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
Copy link
Contributor

@MikeDvorskiy MikeDvorskiy Dec 27, 2024

Choose a reason for hiding this comment

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

  1. 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).

  2. 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?

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Projects
None yet
Development

Successfully merging this pull request may close these issues.

2 participants