[SYCL] Align sub-group implementation and docs #1922

Pennycook · 2020-06-18T17:32:28Z

Deprecate two-input shuffles
Fix linear_id type (size_t => uint32_t)
Add linear ID queries
Remove fence_space from barrier

Signed-off-by: John Pennycook [email protected]

Brings the sub-group implementation closer in line with the specification update from #1565. Includes some minor refactoring to simplify the implementation of requested features (e.g. #1885 and #1916).

- Remove previously deprecated functions - Deprecate two-input shuffles - Fix linear_id type (size_t => uint32_t) Signed-off-by: John Pennycook <[email protected]>

Preparation for updated sub-group features. Combining files to avoid accidental interface differences. Signed-off-by: John Pennycook <[email protected]>

Signed-off-by: John Pennycook <[email protected]>

- get_group_range() changed from uint32_t to range<1> - get_uniform_group_range() no longer supported Signed-off-by: John Pennycook <[email protected]>

Signed-off-by: John Pennycook <[email protected]>

Pennycook · 2020-06-18T22:34:56Z

@bader Could you please take a look at the failing tests for me? I can't see how the failure relates to these changes.

bader · 2020-06-19T09:48:33Z

@bader Could you please take a look at the failing tests for me? I can't see how the failure relates to these changes.

I think these failures are caused by your changes:

D:\BuildBot\Product_worker_intel\sycl-win-x64-pr\llvm.src\sycl\include\CL/sycl/intel/sub_group.hpp(281): error C2491: 'cl::sycl::intel::sub_group::shuffle': definition of dllimport function not allowed
D:\BuildBot\Product_worker_intel\sycl-win-x64-pr\llvm.src\sycl\include\CL/sycl/intel/sub_group.hpp(295): error C2491: 'cl::sycl::intel::sub_group::shuffle_down': definition of dllimport function not allowed
D:\BuildBot\Product_worker_intel\sycl-win-x64-pr\llvm.src\sycl\include\CL/sycl/intel/sub_group.hpp(309): error C2491: 'cl::sycl::intel::sub_group::shuffle_up': definition of dllimport function not allowed

If you think internal testing failures are not related to your patch, please, file a bug report.

Pennycook · 2020-06-19T13:24:14Z

I think these failures are caused by your changes:

Sorry, you're right. I was focusing on the traceback and somehow missed those three lines. Thanks for the help!

Signed-off-by: John Pennycook <[email protected]>

Pennycook · 2020-06-19T13:46:37Z

@alexbatashev : The failures seem related to your change in 1280554, which replaced __SYCL_DEPRECATED__ with __SYCL_EXPORT_DEPRECATED. This patch deprecates three new functions, which correspond to the three build failures on Windows.

Have I made a mistake here, or is there something wrong with using __SYCL_EXPORT_DEPRECATED in this header? I'm not very familiar with compiling for Windows, so I'm not sure why defining these three functions as dllimport is problematic.

alexey-bataev · 2020-06-19T13:49:51Z

@alexey-bataev: The failures seem related to your change in 1280554, which replaced __SYCL_DEPRECATED__ with __SYCL_EXPORT_DEPRECATED. This patch deprecates three new functions, which correspond to the three build failures on Windows.

Have I made a mistake here, or is there something wrong with using __SYCL_EXPORT_DEPRECATED in this header? I'm not very familiar with compiling for Windows, so I'm not sure why defining these three functions as dllimport is problematic.

The author is Alexander Batashev, not Alexey Bataev.

bader · 2020-06-19T13:51:28Z

@alexey-bataev: The failures seem related to your change in 1280554, which replaced __SYCL_DEPRECATED__ with __SYCL_EXPORT_DEPRECATED. This patch deprecates three new functions, which correspond to the three build failures on Windows.
Have I made a mistake here, or is there something wrong with using __SYCL_EXPORT_DEPRECATED in this header? I'm not very familiar with compiling for Windows, so I'm not sure why defining these three functions as dllimport is problematic.

The author is Alexander Batashev, not Alexey Bataev.

I've updated the account id to @alexbatashev in the original comment. Sorry for inconvenience.

Pennycook · 2020-06-19T14:54:52Z

Sorry about the mix-up.

jbrodman · 2020-06-19T17:33:02Z

The comp fails seem like we need to bring back __SYCL_DEPRECATED.

__SYCL_EXPORT_DEPRECATED adding __dllimport is breaking things.

I'm fine with going ahead with that.

__SYCL_EXPORT_DEPRECATED() should not apply to header-only functions. Signed-off-by: John Pennycook <[email protected]>

Pennycook · 2020-06-19T17:54:45Z

The comp fails seem like we need to bring back __SYCL_DEPRECATED.

Thanks, @jbrodman. I've reintroduced __SYCL_DEPRECATED in 90d5156.

alexbatashev · 2020-06-20T06:44:34Z

@alexbatashev : The failures seem related to your change in 1280554, which replaced __SYCL_DEPRECATED__ with __SYCL_EXPORT_DEPRECATED. This patch deprecates three new functions, which correspond to the three build failures on Windows.

Have I made a mistake here, or is there something wrong with using __SYCL_EXPORT_DEPRECATED in this header? I'm not very familiar with compiling for Windows, so I'm not sure why defining these three functions as dllimport is problematic.

The original __SYCL_DEPRECATED__ was removed, because there was no use for it. @jbrodman's answer is correct.

AlexeySachkov

LGTM

Meaning of barrier() was ambiguous, and could have meant: - New barrier() with no fence_space argument - Old barrier() with default fence_space argument Signed-off-by: John Pennycook <[email protected]>

vladimirlaz · 2020-06-26T12:20:23Z

please resolve conflict.

Signed-off-by: John Pennycook <[email protected]>

vladimirlaz · 2020-06-27T17:54:23Z

One more conflict in place...

Signed-off-by: John Pennycook <[email protected]>

Pennycook · 2020-06-29T17:00:03Z

Removing the sub-group collectives is leading to CI failures, so I've put them back in for now to help us close this PR. We should revisit and remove the deprecated functions later.

Signed-off-by: John Pennycook <[email protected]>

Pennycook · 2020-06-29T17:41:06Z

Really sorry @AlexeySachkov; the previous commit didn't include all of my local changes, and I didn't catch it until buildbot failed. To save you from having to review this over and over, I'll ping you when the tests have passed.

Pennycook · 2020-06-30T00:52:29Z

@AlexeySachkov, @vladimirlaz: The tests all passed this time, so I think this is ready for merge. Thanks for your help with this.

Optimized LLVM IR started to contain this intrinsic recently after InstCombine updates, like: "Fold and/or of fcmp into class" ( 08f0388 ) There is no direct counterpart for it in SPIR-V, so testing is mapped on a sequence of SPIR-V instructions: test for both qnan and snan -> OpIsNan test for either of qnan and snan -> isquiet(V) ==> abs(V) >= (unsigned(Inf) | quiet_bit) issignaling(V) ==> isnan(V) && !isquiet(V) test for neg/pos inf -> OpIsInf + OpSignBitSet test for neg/pos normal -> OpIsNormal + OpSignBitSet test for neg/pos subnormal -> issubnormal(V) ==> unsigned(abs(V) - 1) < (all mantissa bits set) test for neg/pos zero -> compare bitcasted to int float with 0 + OpSignBitSet Signed-off-by: Sidorov, Dmitry <[email protected]> Original commit: KhronosGroup/SPIRV-LLVM-Translator@0617edd

Pennycook added 6 commits June 18, 2020 08:03

[SYCL] Align sub-group implementation and docs

369d9dd

- Remove previously deprecated functions - Deprecate two-input shuffles - Fix linear_id type (size_t => uint32_t) Signed-off-by: John Pennycook <[email protected]>

[SYCL] Combine device and host sub-group headers

9788887

Preparation for updated sub-group features. Combining files to avoid accidental interface differences. Signed-off-by: John Pennycook <[email protected]>

[SYCL] Add linear ID queries to sub-group class

8ddc993

Signed-off-by: John Pennycook <[email protected]>

[SYCL] Remove fence_space from sub-group barrier

86d0210

Signed-off-by: John Pennycook <[email protected]>

[SYCL] Fix syntax in sub-group tests

a8e3417

- get_group_range() changed from uint32_t to range<1> - get_uniform_group_range() no longer supported Signed-off-by: John Pennycook <[email protected]>

[SYCL] Run clang-format-9 on sub_group.hpp

9abb244

Signed-off-by: John Pennycook <[email protected]>

Pennycook added the spec extension All issues/PRs related to extensions specifications label Jun 18, 2020

Pennycook requested a review from AlexeySachkov June 18, 2020 17:32

Pennycook requested a review from a team as a code owner June 18, 2020 17:32

Pennycook assigned bader Jun 18, 2020

Pennycook requested a review from vladimirlaz June 18, 2020 17:32

[SYCL] Extract id component 0 before cast

838001e

Signed-off-by: John Pennycook <[email protected]>

[SYCL] Remove duplicated using statements

b1b18ad

Signed-off-by: John Pennycook <[email protected]>

[SYCL] Reintroduce __SYCL_DEPRECATED macro

90d5156

__SYCL_EXPORT_DEPRECATED() should not apply to header-only functions. Signed-off-by: John Pennycook <[email protected]>

vladimirlaz previously approved these changes Jun 22, 2020

View reviewed changes

vladimirlaz self-requested a review June 22, 2020 07:14

AlexeySachkov previously approved these changes Jun 22, 2020

View reviewed changes

[SYCL] Remove default argument from barrier

048dca0

Meaning of barrier() was ambiguous, and could have meant: - New barrier() with no fence_space argument - Old barrier() with default fence_space argument Signed-off-by: John Pennycook <[email protected]>

Pennycook dismissed stale reviews from AlexeySachkov and vladimirlaz via 048dca0 June 23, 2020 13:47

Merge remote-tracking branch 'origin/sycl' into sub-group-tidy

011139d

Signed-off-by: John Pennycook <[email protected]>

Pennycook added 2 commits June 29, 2020 07:23

Merge remote-tracking branch 'origin/sycl' into sub-group-tidy

ca39fa6

Signed-off-by: John Pennycook <[email protected]>

[SYCL] Restore sub-group collectives

b382ba6

Signed-off-by: John Pennycook <[email protected]>

AlexeySachkov previously approved these changes Jun 29, 2020

View reviewed changes

[SYCL] Add #ifdef __SYCL_DEVICE_ONLY__ guards

dad7d1a

Signed-off-by: John Pennycook <[email protected]>

Pennycook dismissed AlexeySachkov’s stale review via dad7d1a June 29, 2020 17:39

jbrodman requested a review from AlexeySachkov June 30, 2020 01:04

vladimirlaz approved these changes Jun 30, 2020

View reviewed changes

AlexeySachkov approved these changes Jun 30, 2020

View reviewed changes

bader merged commit bea6aa2 into intel:sycl Jun 30, 2020

Pennycook deleted the sub-group-tidy branch June 30, 2020 17:29

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Align sub-group implementation and docs #1922

[SYCL] Align sub-group implementation and docs #1922

Pennycook commented Jun 18, 2020 •

edited

Loading

Pennycook commented Jun 18, 2020

bader commented Jun 19, 2020

Pennycook commented Jun 19, 2020

Pennycook commented Jun 19, 2020 •

edited by bader

Loading

alexey-bataev commented Jun 19, 2020

bader commented Jun 19, 2020

Pennycook commented Jun 19, 2020

jbrodman commented Jun 19, 2020

Pennycook commented Jun 19, 2020

alexbatashev commented Jun 20, 2020

AlexeySachkov left a comment

vladimirlaz commented Jun 26, 2020

vladimirlaz commented Jun 27, 2020

Pennycook commented Jun 29, 2020

Pennycook commented Jun 29, 2020

Pennycook commented Jun 30, 2020

[SYCL] Align sub-group implementation and docs #1922

[SYCL] Align sub-group implementation and docs #1922

Conversation

Pennycook commented Jun 18, 2020 • edited Loading

Pennycook commented Jun 18, 2020

bader commented Jun 19, 2020

Pennycook commented Jun 19, 2020

Pennycook commented Jun 19, 2020 • edited by bader Loading

alexey-bataev commented Jun 19, 2020

bader commented Jun 19, 2020

Pennycook commented Jun 19, 2020

jbrodman commented Jun 19, 2020

Pennycook commented Jun 19, 2020

alexbatashev commented Jun 20, 2020

AlexeySachkov left a comment

Choose a reason for hiding this comment

vladimirlaz commented Jun 26, 2020

vladimirlaz commented Jun 27, 2020

Pennycook commented Jun 29, 2020

Pennycook commented Jun 29, 2020

Pennycook commented Jun 30, 2020

Pennycook commented Jun 18, 2020 •

edited

Loading

Pennycook commented Jun 19, 2020 •

edited by bader

Loading