[SYCL] Initial implementation of `invoke_simd` and `uniform` extensions. #5871

kbobrovs · 2022-03-24T06:39:25Z

invoke_simd:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_invoke_simd.asciidoc
Current limitations:
- bool parameter type not supported
Based on POC by @rolandschulz.
uniform:
https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_uniform.asciidoc

Signed-off-by: Konstantin S Bobrovsky [email protected]

s-kanaev

LGTM

sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp

sycl/include/sycl/ext/oneapi/experimental/uniform.hpp

kbobrovs · 2022-03-29T01:18:09Z

@rolandschulz, please also review

sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp

sycl/test/invoke_simd/invoke_simd.cpp

kbobrovs · 2022-03-31T04:29:29Z

Had to rebase to import some recent fixes in other places.

@rolandschulz, @Pennycook, @s-kanaev - please review

s-kanaev

LGTM

- invoke_simd: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_invoke_simd.asciidoc Current limitations: * function pointers and lambdas not supported * simd width fixed to 16, must be auto-deduced in future * bool parameter type not supported - uniform: https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_uniform.asciidoc Signed-off-by: Konstantin S Bobrovsky <[email protected]>

@rolandschulz

- Improvements based on @rolandschulz's POC: * auto-deduce the subgroup size * support lambdas and functors - Added more compilation test cases - Split __builtin_invoke_simd into two overloads - Added TODO about uniform handling in the compiler

- fix compilation errors of other tests related to incorrect sub_group namespace in forward declaration in uniform.hpp - fix test to add SYCL_EXTERNAL to undefined functor's '()' operators Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

…hulz

…hulz.

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs · 2022-04-05T21:26:28Z

The only failure is unrelated - clang FE crash on Printf/long.cpp:

[2022-04-05T02:26:01.584Z] ******************** TEST 'SYCL :: Printf/long.cpp' FAILED ********************
...
[2022-04-05T02:26:01.584Z] $ "w:/jenkins-dir/workspace/SYCL_CI/intel/Win/Test_Suite/llvm.obj/bin/clang-cl.exe" "/EHsc" "-fsycl" "-fsycl-targets=spir64" "W:\jenkins-dir\workspace\SYCL_CI\intel\Win\Test_Suite\llvm-test-suite\SYCL\Printf\long.cpp" "-o" "W:\jenkins-dir\workspace\SYCL_CI\intel\Win\Test_Suite\build\SYCL\Printf\Output\long.cpp.tmp.out"
[2022-04-05T02:26:01.584Z] # command stderr:
[2022-04-05T02:26:01.584Z] PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace, preprocessed source, and associated run script.
[2022-04-05T02:26:01.584Z] Stack dump:
[2022-04-05T02:26:01.585Z] 0.	Program arguments: w:\\jenkins-dir\\workspace\\sycl_ci\\intel\\win\\test_suite\\llvm.obj\\bin\\clang-cl.exe -cc1 ...
...
[2022-04-05T02:26:01.585Z] clang-cl: error: clang frontend command failed due to signal (use -v to see invocation)

kbobrovs · 2022-04-05T21:32:55Z

@rolandschulz, @Pennycook, @s-kanaev, please take one more look. I had to fight some of Linux-only failures (bug in glibcxx, different assert_is_func instantiations mapping to the same mangling in the test), sub_group namespace, as well as had to rebase to include a fix for incorrect sycl::detail::boost::mp11 deploy directory. I think the patch will be ready to merge, as testing passed except one unrelated failure.

v-klochkov · 2022-04-06T14:50:48Z

sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp

+#else
+{
+  // __builtin_invoke_simd is not supported on the host device yet
+  throw sycl::feature_not_supported();


Please do not use the deprecated API. I recently made corresponding change to fix warnings: #5861
(There are few more similar places below)

Suggested change

throw sycl::feature_not_supported();

__ESIMD_UNSUPPORTED_ON_HOST;

thanks, will address in a separate PR.

sycl/include/sycl/ext/oneapi/experimental/uniform.hpp

Pennycook

Approving because I don't want to block this. But I suspect that we will have to revisit the sub_group question I raised in #5871 (comment).

v-klochkov

Please fix the usage of deprecated feature_not_supported() in follow-up PR.

kbobrovs · 2022-04-06T19:00:02Z

Approving because I don't want to block this. But I suspect that we will have to revisit the sub_group question I raised in #5871 (comment).

John, does this address the concern - #5871 (comment) ?
I'm merging since you and others approved, will follow-up on this later if needed.

The uniform extension was implemented in intel#5871. Move the extension specification document to the "experimental" directory.

The uniform extension was implemented in #5871. Move the extension specification document to the "experimental" directory, and update to the latest specification template. Also fix a small bug in the implementation. The `sycl::id` type should not be disallowed for use in `uniform`. Although, it is invalid to pass the work-item's ID to `uniform` (because that value is not uniform across all work-items), the user could construct their own `id` that does have a uniform value.

kbobrovs requested review from a team as code owners March 24, 2022 06:39

kbobrovs requested review from s-kanaev and rolandschulz March 24, 2022 06:39

kbobrovs force-pushed the invoke_simd_api branch from 67891ad to b8ba8ab Compare March 24, 2022 06:42

s-kanaev previously approved these changes Mar 24, 2022

View reviewed changes

sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp Outdated Show resolved Hide resolved

Pennycook reviewed Mar 28, 2022

View reviewed changes

sycl/include/sycl/ext/oneapi/experimental/uniform.hpp Show resolved Hide resolved

kbobrovs dismissed s-kanaev’s stale review via 93c6c39 March 29, 2022 01:17

kbobrovs requested review from Pennycook and s-kanaev March 29, 2022 01:17

Pennycook reviewed Mar 29, 2022

View reviewed changes

sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp Outdated Show resolved Hide resolved

sycl/test/invoke_simd/invoke_simd.cpp Show resolved Hide resolved

kbobrovs force-pushed the invoke_simd_api branch from 6d0c3f7 to 91ce864 Compare March 31, 2022 04:25

s-kanaev requested a review from Pennycook March 31, 2022 11:08

s-kanaev previously approved these changes Mar 31, 2022

View reviewed changes

Pennycook previously approved these changes Mar 31, 2022

View reviewed changes

kbobrovs dismissed stale reviews from Pennycook and s-kanaev via 9d1e5c3 April 1, 2022 02:07

kbobrovs force-pushed the invoke_simd_api branch from 9d1e5c3 to 3df8dad Compare April 1, 2022 02:08

kbobrovs marked this pull request as draft April 1, 2022 18:44

kbobrovs added 9 commits April 4, 2022 08:41

Fix compilation error on some tests

750309c

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

replace sycl::ext::oneapi::sub_group -> sycl::sub_group

2b9694e

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

Fix sub_group-related compilation errors, and a test.

bcda0ee

W/a __regcall function std::is_function_v problem on Linux.

fddd978

Added is_function_v problem root cause description found by @rolandsc…

16e6ec7

…hulz.

Formalize is_function_v w/a, extend invoke_simd test.

b910f4f

Signed-off-by: Konstantin S Bobrovsky <[email protected]>

kbobrovs force-pushed the invoke_simd_api branch from a767c1c to b910f4f Compare April 4, 2022 18:18

Fix Linux compilation errors.

83c6872

kbobrovs force-pushed the invoke_simd_api branch from de3a83b to 83c6872 Compare April 4, 2022 23:59

kbobrovs added a commit to kbobrovs/llvm that referenced this pull request Apr 5, 2022

Update __builtin_invoke_simd mangling according to PR intel#5871

92a9b9a

kbobrovs mentioned this pull request Apr 5, 2022

[SYCL] Add LowerInvokeSimd LLVM pass to lower invoke_simd for BEs' consumption. #5864

Merged

kbobrovs marked this pull request as ready for review April 5, 2022 21:26

s-kanaev approved these changes Apr 6, 2022

View reviewed changes

v-klochkov reviewed Apr 6, 2022

View reviewed changes

Pennycook reviewed Apr 6, 2022

View reviewed changes

sycl/include/sycl/ext/oneapi/experimental/uniform.hpp Show resolved Hide resolved

kbobrovs requested a review from Pennycook April 6, 2022 15:44

Pennycook approved these changes Apr 6, 2022

View reviewed changes

v-klochkov approved these changes Apr 6, 2022

View reviewed changes

kbobrovs merged commit a37ca84 into intel:sycl Apr 6, 2022

kbobrovs deleted the invoke_simd_api branch April 6, 2022 19:59

kbobrovs mentioned this pull request Apr 6, 2022

[SYCL] Replace deprecated sycl::feature_not_supported in invoke_simd.hpp #5976

Merged

gmlueck added a commit to gmlueck/llvm that referenced this pull request Sep 21, 2023

[SYCL][Doc] Move uniform extension spec

c8b74f2

The uniform extension was implemented in intel#5871. Move the extension specification document to the "experimental" directory.

gmlueck mentioned this pull request Sep 21, 2023

[SYCL][Doc] Move uniform extension spec #11259

Merged

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[SYCL] Initial implementation of `invoke_simd` and `uniform` extensions. #5871

[SYCL] Initial implementation of `invoke_simd` and `uniform` extensions. #5871

Uh oh!

kbobrovs commented Mar 24, 2022 •

edited

Loading

Uh oh!

s-kanaev left a comment

Uh oh!

Uh oh!

Uh oh!

kbobrovs commented Mar 29, 2022

Uh oh!

Uh oh!

Uh oh!

kbobrovs commented Mar 31, 2022

Uh oh!

s-kanaev left a comment

Uh oh!

kbobrovs commented Apr 5, 2022

Uh oh!

kbobrovs commented Apr 5, 2022

Uh oh!

v-klochkov Apr 6, 2022 •

edited

Loading

Uh oh!

v-klochkov Apr 6, 2022

Uh oh!

kbobrovs Apr 6, 2022

Uh oh!

Uh oh!

Pennycook left a comment

Uh oh!

v-klochkov left a comment

Uh oh!

kbobrovs commented Apr 6, 2022 •

edited

Loading

Uh oh!

Uh oh!

	throw sycl::feature_not_supported();
	__ESIMD_UNSUPPORTED_ON_HOST;

[SYCL] Initial implementation of invoke_simd and uniform extensions. #5871

[SYCL] Initial implementation of invoke_simd and uniform extensions. #5871

Uh oh!

Conversation

kbobrovs commented Mar 24, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

s-kanaev left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

kbobrovs commented Mar 29, 2022

Uh oh!

Uh oh!

Uh oh!

kbobrovs commented Mar 31, 2022

Uh oh!

s-kanaev left a comment

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented Apr 5, 2022

Uh oh!

kbobrovs commented Apr 5, 2022

Uh oh!

v-klochkov Apr 6, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

v-klochkov Apr 6, 2022

Choose a reason for hiding this comment

Uh oh!

kbobrovs Apr 6, 2022

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Pennycook left a comment

Choose a reason for hiding this comment

Uh oh!

v-klochkov left a comment

Choose a reason for hiding this comment

Uh oh!

kbobrovs commented Apr 6, 2022 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

[SYCL] Initial implementation of `invoke_simd` and `uniform` extensions. #5871

[SYCL] Initial implementation of `invoke_simd` and `uniform` extensions. #5871

kbobrovs commented Mar 24, 2022 •

edited

Loading

v-klochkov Apr 6, 2022 •

edited

Loading

kbobrovs commented Apr 6, 2022 •

edited

Loading