Fall back to SYCL 2020 API for a generic SYCL implementation #1954

dmitriy-sobolev · 2024-12-03T16:19:39Z

Let's assume that a generic SYCL library is compliant to SYCL 2020 and fallback to that API.

Below is a specific example of the issue with sycl::property::no_init. Other places fixed in the PR are essentially the same.

Before:

#define _ONEDPL_NO_INIT_PRESENT (_ONEDPL_LIBSYCL_VERSION >= 50300)
...
using __no_init =
#if _ONEDPL_NO_INIT_PRESENT
    // SYCL 2020 API; Will be selected only with newer DPC++ versions
    sycl::property::no_init;
#else
    // Specific to older DPC++; Will be selected with older DPC++ and other SYCL implementations
    sycl::property::noinit; 
#endif

After:

#define _ONEDPL_LIBSYCL_VERSION_LESS_THAN(_ONEDPL_SYCL_VERSION_NUMBER)                                                 \
    (_ONEDPL_LIBSYCL_VERSION && _ONEDPL_LIBSYCL_VERSION < _ONEDPL_SYCL_VERSION_NUMBER)
#define _ONEDPL_SYCL2020_NO_INIT_PRESENT (!_ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300))
...
using __no_init =
#if _ONEDPL_SYCL2020_NO_INIT_PRESENT
    // SYCL 2020 API; Will be selected with newer DPC++ versions and other SYCL implementations
    sycl::property::no_init;
#elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300)
    // Specific to older DPC++; Will be selected only with older DPC++
    sycl::property::noinit;
#endif

Meanwhile, the conditions relying on DPC++ extensions remained the same, e.g. FPGA support or ESIMD.

Additional changes:

Some macros were renamed to have SYCL2020 or LIBSYCL in their names to denote their purpose (either SYCL 2020 feature or dependency on DPC++ SYCL implementation/extension).

_ONEDPL_SYCL2020_UNIFIED_USM_BUFFER_PRESENT was renamed to _ONEDPL_SYCL2020_DEFAULT_ACCESSOR_CONSTRUCTOR_PRESENT to show the exact SYCL 2020 feature needed to enable code section. _ONEDPL_SYCL_L0_EXT_PRESENT was introduced for getting L0 backend with different SYCL implementations.

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort.h

dmitriy-sobolev · 2024-12-09T09:37:10Z

Turning into draft to make the approach more scalable.

akukanov

The approach looks good to me in principle, but I have not checked thoroughly that all the new definitions match the old ones.

test/general/test_policies.pass.cpp

include/oneapi/dpl/pstl/hetero/dpcpp/unseq_backend_sycl.h

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

dmitriy-sobolev · 2024-12-13T10:42:39Z

@MikeDvorskiy could you also take a look at the PR? The macros can be tricky and prone to subtle issues, so the second reviewer can help a lot.

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

MikeDvorskiy · 2024-12-17T17:11:11Z

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

 #    define _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE(SIZE) sycl::reqd_sub_group_size(SIZE)
-#else
+#elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300)


What if both macros are evaluated to zero? (Here and in the others similar places below)

If it is possible, we don't control the all cases and get an error - undefined _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE...

If it is not possible, I don't understand why do we write the second if - #elif...

Of course we do not control all the cases, because we cannot predict what an arbitrary SYCL implementation can or cannot do, and what workaround it would require if something is not supported.

The optimistic expectation is that fully conformant SYCL implementations go into the "present" branch with no effort from our side. However if for some implementation that's not true, the case (1) will happen and we will need to analyze the failure and provide a proper workaround, likely in another #elif branch.

Specifically for _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE, we can define it as no-op in the #else branch. But this might result in runtime errors for oneDPL kernels that really require a certain WG size, which would be even worse than the guaranteed compilation error.

Another approach is to switch from "present" macros to "workaround" macros. I think Dmitry tried something of that kind with "absent" macros in a previous revision of the patch. It would be something like this:

#if _ONEDPL_LIBSYCL_REQD_SUB_GROUP_SIZE_WA // == _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300) # define _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE(SIZE) intel::reqd_sub_group_size(SIZE) #else // SYCL 2020 conformance expected # define _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE(SIZE) sycl::reqd_sub_group_size(SIZE) #endif

For an arbitrary implementation of SYCL it is no better than the proposed way, because if sycl::reqd_sub_group_size is not supported, there will also be a compilation error. And while the code appears to be slightly simpler, a downside is that the conformant code comes last, not first.

Let me add more details on top of Alexey's answer.

What if both macros are evaluated to zero... If it is possible

It is possible.

Let's consider such a case: older versions of FOOSYCL, some SYCL library, do not support sycl::reqd_sub_group_size. There is no workaround, or it was simply forgotten to be added.

#define _ONEDPL_SYCL2020_REQD_SUB_GROUP_SIZE_PRESENT \ (!_ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300) && !_ONEDPL_FOOSYCL_VERSION_LESS_THAN(42)) // Current approach: older versions of FOOSYCL will have undefined macro, which may have unpredictable consequences #if _ONEDPL_SYCL2020_REQD_SUB_GROUP_SIZE_PRESENT # define _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE(SIZE) sycl::reqd_sub_group_size(SIZE) #elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300) # define _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE(SIZE) intel::reqd_sub_group_size(SIZE) #endif // Possible approach: error will be emitted with the older FOOSYCL versions, but it is more verbose #if _ONEDPL_SYCL2020_REQD_SUB_GROUP_SIZE_PRESENT # define _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE(SIZE) sycl::reqd_sub_group_size(SIZE) #elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300) # define _ONEDPL_SYCL_REQD_SUB_GROUP_SIZE(SIZE) intel::reqd_sub_group_size(SIZE) #else # error "No implementation for reqd_sub_group_size" #endif

Was it your concern, Mikhail? If so, would you prefer "Possible approach" above? It's an edge-case, I am reluctant to do that across the PR due to the verbosity.

we can define it as no-op in the #else branch

We can also define #pragma error in the #else branch...
I think compile time error is better then runtime error.

if sycl::reqd_sub_group_size is not supported, there will also be a compilation error

Ok, this behavior is also acceptable for me.

# error "No implementation for reqd_sub_group_size"

Yes, this approach is acceptable to me

It's an edge-case, I am reluctant to do that across the PR due to the verbosity.

Sorry, but it seems I didn't catch what you want to say...

We can also define #pragma error in the #else branch...
I think compile time error is better then runtime error.

Done.

It's an edge-case, I am reluctant to do that across the PR due to the verbosity.

It was about proliferation of "#else... #error ..." parts, which I assumed as non-necessary. It is not relevant, I changed my mind: it is better to have these parts to avoid UB in some cases.

Signed-off-by: Dmitriy Sobolev <[email protected]>

…N_LESS_THAN(version)

This reverts commit 4990728.

MikeDvorskiy · 2024-12-20T13:56:00Z

include/oneapi/dpl/pstl/hetero/dpcpp/utils_ranges_sycl.h

 inline constexpr all_host_view_fn
-#else
+#elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(60200)
 inline constexpr all_view_fn<sycl::access::mode::read_write, __dpl_sycl::__host_target,
                             sycl::access::placeholder::false_t>


#else # error ... ?

MikeDvorskiy · 2024-12-20T14:03:38Z

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

    sycl::property::no_init;
-#else
+#elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300)
    sycl::property::noinit;
 #endif


#else #error ..... branch ?

MikeDvorskiy · 2024-12-20T14:03:55Z

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

 using __sub_group = sycl::sub_group;
-#else
+#elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50700)
 using __sub_group = sycl::ONEAPI::sub_group;
 #endif


#else #error ..... branch ?

MikeDvorskiy · 2024-12-20T14:04:01Z

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

    return __buffer.size();
-#else
+#elif _ONEDPL_LIBSYCL_VERSION_LESS_THAN(50300)
    return __buffer.get_count();
 #endif


#else #error ..... branch ?
Here and in the rest same places below.

MikeDvorskiy

LGTM

dmitriy-sobolev requested review from MikeDvorskiy, SergeyKopienko, mmichel11, danhoeflinger and adamfidel December 3, 2024 16:19

dmitriy-sobolev added enhancement bug labels Dec 3, 2024

dmitriy-sobolev changed the title ~~Fall back to SYCL 2020 API for unknown SYCL library~~ Fall back to SYCL 2020 API for a generic SYCL library Dec 3, 2024

dmitriy-sobolev changed the title ~~Fall back to SYCL 2020 API for a generic SYCL library~~ Fall back to SYCL 2020 API for a generic SYCL implementation Dec 3, 2024

dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/sycl2020-fallback branch from c0fa332 to 91d07bc Compare December 3, 2024 17:48

dmitriy-sobolev marked this pull request as ready for review December 4, 2024 11:20

akukanov reviewed Dec 4, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h Outdated Show resolved Hide resolved

akukanov reviewed Dec 4, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort.h Outdated Show resolved Hide resolved

dmitriy-sobolev self-assigned this Dec 9, 2024

dmitriy-sobolev marked this pull request as draft December 9, 2024 09:36

dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/sycl2020-fallback branch 4 times, most recently from 71e5673 to a0fc677 Compare December 9, 2024 13:48

dmitriy-sobolev marked this pull request as ready for review December 10, 2024 10:52

akukanov reviewed Dec 11, 2024

View reviewed changes

test/general/test_policies.pass.cpp Outdated Show resolved Hide resolved

include/oneapi/dpl/pstl/hetero/dpcpp/unseq_backend_sycl.h Outdated Show resolved Hide resolved

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h Outdated Show resolved Hide resolved

dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/sycl2020-fallback branch from f321c2c to 865b88a Compare December 12, 2024 09:33

dmitriy-sobolev added this to the 2022.8.0 milestone Dec 16, 2024

SergeyKopienko reviewed Dec 17, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h Show resolved Hide resolved

SergeyKopienko reviewed Dec 17, 2024

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h Outdated Show resolved Hide resolved

MikeDvorskiy reviewed Dec 17, 2024

View reviewed changes

dmitriy-sobolev added 2 commits December 19, 2024 19:45

Fall back to SYCL 2019 API

e94bc36

Signed-off-by: Dmitriy Sobolev <[email protected]>

Workaround missing _ONEDPL_DPCPP_LIBSYCL_ZERO_OR_GE

40a4bee

Signed-off-by: Dmitriy Sobolev <[email protected]>

dmitriy-sobolev added 18 commits December 19, 2024 19:45

Rename two more macros to align with others

a90825f

Small fix

67dd043

Generalize the approach, introduce _ONEDPL_GENERIC_SYCL_LIBRARY macro

0c3da6f

Make sycl feature selection logic more scalable

b555202

Simplify

47a9bd9

Clarify the comment

ebc771a

Restore _ONEDPL_SYCL2020_KERNEL_DEVICE_API_ABSENT condition

bf16731

Bump up placeholder host accessor condition to the original value

a6d9e6e

absent -> present logic

cd0c556

Correct some versions

2e41072

Replace _ONEDPL_LIBSYCL_VERSION < version with _ONEDPL_LIBSYCL_VERSIO…

e552b17

…N_LESS_THAN(version)

Get rid of _ONEDPL_LIBSYCL_VERSION use in a test

c2e2f2c

Restore the removed test-case, minor stylistic corrections

5adf58f

Get rid of comment cluttering

05f0cec

Use ext_oneapi_level_zero with DPC++ only

bbd6c82

Stylistic changes

68e7eb5

Revert "Stylistic changes"

a84e9c5

This reverts commit 4990728.

Simplified stylistic changes

5a9f7fd

dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/sycl2020-fallback branch from ae90d1e to 5a9f7fd Compare December 19, 2024 19:48

Fix issues after rebase

fab2611

MikeDvorskiy reviewed Dec 20, 2024

View reviewed changes

dmitriy-sobolev added 3 commits December 20, 2024 16:12

Add #else

dbbfd53

Add #else in utils_ranges_sycl.h

3e8c4ab

clang-format

fb15001

MikeDvorskiy approved these changes Dec 20, 2024

View reviewed changes

MikeDvorskiy merged commit c2a226b into main Dec 23, 2024
21 of 22 checks passed

MikeDvorskiy deleted the dev/dmitriy-sobolev/sycl2020-fallback branch December 23, 2024 10:55

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fall back to SYCL 2020 API for a generic SYCL implementation #1954

Fall back to SYCL 2020 API for a generic SYCL implementation #1954

dmitriy-sobolev commented Dec 3, 2024 •

edited

Loading

dmitriy-sobolev commented Dec 9, 2024

akukanov left a comment

dmitriy-sobolev commented Dec 13, 2024

MikeDvorskiy Dec 17, 2024 •

edited

Loading

akukanov Dec 17, 2024 •

edited

Loading

akukanov Dec 17, 2024 •

edited

Loading

akukanov Dec 17, 2024

dmitriy-sobolev Dec 18, 2024 •

edited

Loading

MikeDvorskiy Dec 20, 2024

MikeDvorskiy Dec 20, 2024 •

edited

Loading

MikeDvorskiy Dec 20, 2024

MikeDvorskiy Dec 20, 2024 •

edited

Loading

dmitriy-sobolev Dec 20, 2024

MikeDvorskiy Dec 20, 2024

MikeDvorskiy Dec 20, 2024

MikeDvorskiy Dec 20, 2024

dmitriy-sobolev Dec 20, 2024

MikeDvorskiy Dec 20, 2024 •

edited

Loading

dmitriy-sobolev Dec 20, 2024

MikeDvorskiy left a comment

Fall back to SYCL 2020 API for a generic SYCL implementation #1954

Fall back to SYCL 2020 API for a generic SYCL implementation #1954

Conversation

dmitriy-sobolev commented Dec 3, 2024 • edited Loading

dmitriy-sobolev commented Dec 9, 2024

akukanov left a comment

Choose a reason for hiding this comment

dmitriy-sobolev commented Dec 13, 2024

MikeDvorskiy Dec 17, 2024 • edited Loading

Choose a reason for hiding this comment

akukanov Dec 17, 2024 • edited Loading

Choose a reason for hiding this comment

akukanov Dec 17, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

dmitriy-sobolev Dec 18, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MikeDvorskiy Dec 20, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MikeDvorskiy Dec 20, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MikeDvorskiy Dec 20, 2024 • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

MikeDvorskiy left a comment

Choose a reason for hiding this comment

dmitriy-sobolev commented Dec 3, 2024 •

edited

Loading

MikeDvorskiy Dec 17, 2024 •

edited

Loading

akukanov Dec 17, 2024 •

edited

Loading

akukanov Dec 17, 2024 •

edited

Loading

dmitriy-sobolev Dec 18, 2024 •

edited

Loading

MikeDvorskiy Dec 20, 2024 •

edited

Loading

MikeDvorskiy Dec 20, 2024 •

edited

Loading

MikeDvorskiy Dec 20, 2024 •

edited

Loading