[SYCL][ABI Break] Add support for per-kernel auto GRF mode specification, and reimplement feature using kernel properties #9258

sarnex · 2023-04-28T18:55:45Z

This change reimplements the GRF mode specification using native SYCL kernel properties, does an ABI break, and adds the automatic option.

We add a new property named sycl::detail::register_alloc_mode, and it takes in an enum sycl::detail::register_alloc_mode_enum which currently has two values: automatic and large.

This can be applied to kernels as below:

    properties prop{register_alloc_mode<register_alloc_mode_enum::automatic>};
...
 cgh.parallel_for<class Foo>(
          Size, prop, [=](id<1> i) { PA[i] += 2; });

We do some register_alloc_mode specific work in CompileTimePropertiesPass, we add function metadata named RegisterAllocMode based on the function attribute sycl-register-alloc-mode. This is because llvm-spirv looks for the metadata. This metadata is how AOT works.

We also do some register_alloc_mode specific work in sycl-post-link:

Split based on the value of the sycl-register-alloc-mode function attribute added in the front end
Add a binary property named sycl-register-alloc-mode used in the SYCL runtime.

Note that ESIMD does not work at all yet for automatic yet, not for JIT nor AOT. This is because the VC backend does not support auto GRF. I have made a feature request for the VC team to add this. large continues to work for JIT only.

llvm/tools/sycl-post-link/sycl-post-link.cpp

llvm/tools/sycl-post-link/ModuleSplitter.cpp

sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp

sycl/source/detail/program_manager/program_manager.cpp

llvm/test/tools/sycl-post-link/sycl-auto-grf.ll

llvm/test/SYCLLowerIR/lower_kernel_props.ll

llvm/lib/SYCLLowerIR/LowerKernelProps.cpp

sarnex · 2023-05-04T15:05:08Z

@AlexeySachkov I've reimplemented this using kernel properties and did the ABI-breaking device binary property rename. Please take a look when you get a chance, thanks.

AlexeySachkov

It is always good to see the code removed! However, I'm concerned that lowering of the property is done by sycl-post-link and not by a dedicated pass

llvm/tools/sycl-post-link/sycl-post-link.cpp

sarnex · 2023-05-05T19:30:49Z

However, I'm concerned that lowering of the property is done by sycl-post-link and not by a dedicated pass

fixed in latest commit, thx

AlexeySachkov

sycl-post-link and SYCLLowerIR parts LGTM. Thanks for the refactoring and unification!

llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-alloc-mode.ll

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp

sarnex · 2023-05-09T13:38:45Z

@intel/llvm-reviewers-runtime @intel/dpcpp-esimd-reviewers This one is ready for review now, the major feedback has been addressed. Thanks!

sarnex · 2023-05-12T13:56:53Z

sorry for the force push, the e2e test changes made rebasing preserving commits difficult

sarnex · 2023-05-15T13:35:26Z

@intel/llvm-reviewers-runtime ping on this one, thanks

sarnex · 2023-05-16T14:02:30Z

@bso-intel ping, thanks.

steffenlarsen

Since this is marked as an ABI break, we need a bump to SYCL_DEV_ABI_VERSION in accordance with the ABI policy.

Also, I feel like I got the answer with the previous implementation of GRF mode handling, but is this new property intended for users? If yes, I think we need an extension documenting it. If not, we should consider placing it in another namespace to avoid confusion.

sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp

sycl/source/detail/program_manager/program_manager.cpp

sycl/include/sycl/ext/intel/experimental/kernel_properties.hpp

sycl/source/detail/program_manager/program_manager.cpp

…implement feature using kernel properties This works extends the existing support we have for large GRF mode specification. I introduce a new argument to set_kernel_properties, kernel_properties::use_auto_grf. Then, we update LowerKernelProps to lower this new attribute and do sanity checks. Next, we update sycl-post-link to split and add an image property based on this property. Finally, we update program manager to check the image property and pass the correct flag for JIT. For AOT, this works through the RegisterAllocMode metadata that we add during LowerKernelProps and is kept through llvm-spirv. Note that ESIMD does not work at all yet, not for JIT nor AOT. This is because the VC backend does not support auto GRF. I have made a feature request for the VC team to add this. Signed-off-by: Sarnie, Nick <[email protected]>

Signed-off-by: Sarnie, Nick <[email protected]>

sarnex · 2023-05-16T17:35:34Z

sorry for rebase, there was a conflict

@steffenlarsen @dm-vodopyanov all feedback should be addressed (hopefully)

steffenlarsen

Runtime changes LGTM!

Signed-off-by: Sarnie, Nick <[email protected]>

dm-vodopyanov

Thanks, RT part LGTM

v-klochkov

Ok for ESIMD. The changes look great. Thank you.
The fix is ABI Breaking and thus may need special approval at PXT meeting.

aelovikov-intel · 2023-05-16T20:36:56Z

We are in ABI breaking window.

sarnex · 2023-05-16T20:38:07Z

@aelovikov-intel we think we may still need approval even within the window. are you sure that we do not?

sarnex · 2023-05-17T13:34:25Z

@aelovikov-intel We confirmed internally we do NOT need approval, it was a false alarm.

@intel/llvm-gatekeepers This one is ready for merge. Thanks!

v-klochkov · 2023-06-05T18:45:11Z

sycl/test-e2e/DeviceCodeSplit/grf.cpp

 #include <sycl/sycl.hpp>

 using namespace sycl;
+using namespace sycl::detail;


sycl::detail::register_alloc_mode ?!

Sorry for the late comment, but I just did not realize this issue during code-review phase and discovered it only now when tried to use the new API.

I believe we must never inspire users to use 'detail' namespace, and ask include files from "*/detail/*" - it is only for SYCL internal needs. Only SYCL implementation should use detail namespace, not user.

The enum must have been defined not as "sycl::detail::register_alloc_mode", but as "sycl::ext::oneapi::experimental::register_alloc_mode".

The file kernel_properties.hpp also needed to be placed at ext/oneapi/experimental/kernel_properties.hpp.

Maybe the file needs to be placed in ext/intel/experimental/kernel_properties.cpp (where it was), which seems more correct than where it is now (sycl/detail).

The detail namespace was picked after the discussion in #9258 (comment) as it was believed to be an internal-only property. If it is meant to be used externally we would normally require documentation of the feature prior to implementation. When we have that, I see no problem in moving it out of the detail namespace.

The previous "experimental" extension should not have been removed while there is no new equivalent "experimental" feature. This means that we need to switch from the old "experimental" to new "detail" (code breaks), then we need to switch from "detail" to "experimental" (another code break), and then finally from "experimental" to final (3rd code break). Is there any reason apart from the lack of documentation to add this to "detail" first?

sarnex temporarily deployed to aws April 28, 2023 19:21 — with GitHub Actions Inactive

sarnex temporarily deployed to aws April 28, 2023 19:55 — with GitHub Actions Inactive

sarnex commented Apr 28, 2023

View reviewed changes

llvm/tools/sycl-post-link/sycl-post-link.cpp Outdated Show resolved Hide resolved

sarnex marked this pull request as ready for review April 28, 2023 20:17

sarnex requested review from a team as code owners April 28, 2023 20:17

sarnex requested a review from bso-intel April 28, 2023 20:17

AlexeySachkov reviewed May 2, 2023

View reviewed changes

sarnex changed the title ~~[SYCL] Add support for per-kernel auto GRF mode specification~~ [SYCL][ABI Break] Add support for per-kernel auto GRF mode specification, and reimplement feature using kernel properties May 4, 2023

sarnex requested a review from a team as a code owner May 4, 2023 15:03

sarnex requested a review from AlexeySachkov May 4, 2023 15:05

sarnex force-pushed the autogrf branch from d257c58 to 623bcd3 Compare May 4, 2023 15:11

sarnex temporarily deployed to aws May 4, 2023 19:03 — with GitHub Actions Inactive

sarnex temporarily deployed to aws May 5, 2023 01:04 — with GitHub Actions Inactive

AlexeySachkov reviewed May 5, 2023

View reviewed changes

llvm/tools/sycl-post-link/sycl-post-link.cpp Outdated Show resolved Hide resolved

llvm/tools/sycl-post-link/sycl-post-link.cpp Outdated Show resolved Hide resolved

sarnex temporarily deployed to aws May 5, 2023 20:38 — with GitHub Actions Inactive

sarnex temporarily deployed to aws May 6, 2023 02:50 — with GitHub Actions Inactive

AlexeySachkov approved these changes May 9, 2023

View reviewed changes

llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/register-alloc-mode.ll Show resolved Hide resolved

llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp Show resolved Hide resolved

sarnex temporarily deployed to aws May 9, 2023 13:41 — with GitHub Actions Inactive

sarnex force-pushed the autogrf branch from b7cc85c to 829c1ca Compare May 9, 2023 13:46

sarnex mentioned this pull request May 9, 2023

Refactor argument order of 'attributeToExecModeMetadata' in CompileTimePropertiesPass #9370

Closed

sarnex temporarily deployed to aws May 9, 2023 14:13 — with GitHub Actions Inactive

sarnex temporarily deployed to aws May 9, 2023 14:53 — with GitHub Actions Inactive

sarnex force-pushed the autogrf branch from 829c1ca to e219d6c Compare May 12, 2023 13:56

sarnex temporarily deployed to aws May 12, 2023 14:29 — with GitHub Actions Inactive

sarnex temporarily deployed to aws May 12, 2023 15:14 — with GitHub Actions Inactive

steffenlarsen requested changes May 16, 2023

View reviewed changes

dm-vodopyanov reviewed May 16, 2023

View reviewed changes

sarnex temporarily deployed to aws May 16, 2023 16:41 — with GitHub Actions Inactive

sarnex added 4 commits May 16, 2023 10:27

address runtime fedback

acc7dd2

Signed-off-by: Sarnie, Nick <[email protected]>

whitespace

942ff8c

Signed-off-by: Sarnie, Nick <[email protected]>

use detail namespace

96b292d

Signed-off-by: Sarnie, Nick <[email protected]>

sarnex force-pushed the autogrf branch from b0a113c to 96b292d Compare May 16, 2023 17:32

steffenlarsen approved these changes May 16, 2023

View reviewed changes

sarnex temporarily deployed to aws May 16, 2023 17:58 — with GitHub Actions Inactive

fix lit test

0865098

Signed-off-by: Sarnie, Nick <[email protected]>

sarnex temporarily deployed to aws May 16, 2023 18:42 — with GitHub Actions Inactive

dm-vodopyanov approved these changes May 16, 2023

View reviewed changes

sarnex temporarily deployed to aws May 16, 2023 19:14 — with GitHub Actions Inactive

v-klochkov reviewed May 16, 2023

View reviewed changes

v-klochkov approved these changes May 17, 2023

View reviewed changes

dm-vodopyanov merged commit f363bb2 into intel:sycl May 17, 2023

v-klochkov reviewed Jun 5, 2023

View reviewed changes

[SYCL][ABI Break] Add support for per-kernel auto GRF mode specification, and reimplement feature using kernel properties #9258

[SYCL][ABI Break] Add support for per-kernel auto GRF mode specification, and reimplement feature using kernel properties #9258

Uh oh!

Conversation

sarnex commented Apr 28, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

sarnex commented May 4, 2023

Uh oh!

AlexeySachkov left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

sarnex commented May 5, 2023

Uh oh!

AlexeySachkov left a comment • 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!

Uh oh!

Uh oh!

sarnex commented May 9, 2023

Uh oh!

sarnex commented May 12, 2023

Uh oh!

sarnex commented May 15, 2023

Uh oh!

sarnex commented May 16, 2023

Uh oh!

steffenlarsen left a comment

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

sarnex commented May 16, 2023

Uh oh!

steffenlarsen left a comment

Choose a reason for hiding this comment

Uh oh!

dm-vodopyanov left a comment

Choose a reason for hiding this comment

Uh oh!

v-klochkov left a comment • 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!

aelovikov-intel commented May 16, 2023

Uh oh!

sarnex commented May 16, 2023

Uh oh!

sarnex commented May 17, 2023

Uh oh!

v-klochkov Jun 5, 2023

Choose a reason for hiding this comment

Uh oh!

v-klochkov Jun 5, 2023

Choose a reason for hiding this comment

Uh oh!

steffenlarsen Jun 6, 2023

Choose a reason for hiding this comment

Uh oh!

sarnex commented Apr 28, 2023 •

edited

Loading

AlexeySachkov left a comment •

edited

Loading

v-klochkov left a comment •

edited

Loading