[SYCL] Implement initial parts of sycl_ext_oneapi_kernel_properties #6941

steffenlarsen · 2022-10-03T16:24:13Z

This commit implements the following parts of the
sycl_ext_oneapi_kernel_properties extension:

The work_group_size, work_group_size_hint, and sub_group_size properties.
The new overloads for single_task, parallel_for, and parallel_for_work_group in handler.
The new shortcuts for single_task and parallel_for in queue.
Support for merging property lists which is used when kernel functors have a get member for properties. The changes do not include any changes related to the device_has property and corresponding interfaces.

Changes to the Clang frontend:

Fixed an issue preventing initializer lists of characters to be used as strings in add_ir_{attributes|annotations}_* attributes.
The add_ir_attributes_function will now be copied onto the generated kernel declaration from the recognized kernel_* functions.
Added a warning diagnostic for when a kernel has both an add_ir_attributes_function (with values) and one or more potentially conflicting SYCL attributes.

Changes to sycl-post-link:

sycl-post-link is now able to generate metadata for recognized SYCL LLVM IR attributes on functions. These are "sycl-work-group-size", "sycl-work-group-size-hint", and "sycl-sub-group-size".
The previously mentioned new recognized LLVM IR attributes are translated into existing metadata the SPIR-V Translator translates into corresponding SPIR-V execution modes. If these metadata nodes already exist on a function (e.g. added through a SYCL 2020 attribute), the related property is ignored.

Additionally these changes fixes a minor bug with how work-group sizes were not reported in the correct order when invalid work-group sizes are given during launch of a kernel.

This commit implements the following parts of the sycl_ext_oneapi_kernel_properties extension: * The work_group_size, work_group_size_hint, and sub_group_size properties. * The new overloads for single_task, parallel_for, and parallel_for_work_group in handler. * The new shortcuts for single_task and parallel_for in queue. * Support for merging property lists which is used when kernel functors have a get member for properties. The changes do not include any changes related to the device_has property and corresponding interfaces. Changes to the Clang frontend: * Fixed an issue preventing initializer lists of characters to be used as strings in add_ir_{attributes|annotations}_* attributes. * The add_ir_attributes_function will now be copied onto the generated kernel declaration from the recognized kernel_* functions. Changes to sycl-post-link: * sycl-post-link is now able to generate metadata for recognized SYCL LLVM IR attributes on functions. These are "sycl-work-group-size", "sycl-work-group-size-hint", and "sycl-sub-group-size". * The previously mentioned new recognized LLVM IR attributes are translated into existing metadata the SPIR-V Translator translates into corresponding SPIR-V execution modes. If these metadata nodes already exist on a function (e.g. added through a SYCL 2020 attribute), the related property is ignored. Additionally these changes fixes a minor bug with how work-group sizes were not reported in the correct order when invalid work-group sizes are given during launch of a kernel. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

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

sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp

Pennycook · 2022-10-03T18:02:25Z

sycl/include/sycl/handler.hpp

+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;


I just want to check -- what happens if somebody enqueues the same kernel function multiple times, once with properties and once without? I ask because I thought the name was the only connection between a C++ representation of a kernel with the SPIR-V, and I'm wondering if the properties need to be included in the mangling somehow...

That is a very good question! I suspect you might be right that we need to include the properties in the name. I will add some tests for this ASAP!

Based on my testing, it looks like creating two unnamed submissions with the same kernel but different properties (including one with a property and one without) will cause kernel naming conflicts. That is, the compiler will fail because it tries to create two different kernels with the same name.

An option for fixing this is to create a wrapper class to be used by get_kernel_name_t when there are properties. A small proof-of-concept of this is steffenlarsen@b4bc767 but the frontend causes an error in DiagnoseKernelNameType because the kernel isn't unnamed and the instantiation comes from sycl::queue.

@elizabethandrews & @Fznamznon - Do you have some ideas of how we can get around this? Should we somehow force it to consider the wrapper as unnamed too, or would that break other assumptions?

@Pennycook - Do you think this should block this PR or are you okay with having conflicting properties for the same kernel cause name conflicts for now? It is definitely something that should be addressed, but luckily it will be a relaxation rather than a regression in functionality when we get it fixed.

I don't want to block this PR. I want developers to be able to start experimenting with compile-time properties as soon as possible! 😄

Since this is only a partial implementation of something that will be going into the "experimental" namespace initially, I'm not concerned about there being some corner cases that don't work. If it wouldn't be much work, an error that specifically says it's not possible to provide two sets of properties for the same unnamed lambda might save somebody some head-scratching.... But if that's a lot of work, maybe just a FIXME or a TODO would do it.

It might also be worth adding a NOTE to the extension specification about the limitation. This will also remind us to address it before moving the extension out of "experimental".

I have added a note in the extension specification about the current limitation and have added TODO's here and in similar locations.

asudarsa · 2022-10-03T20:03:14Z

Adding @sarnex who is working on a similar support. Thanks

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

clang/include/clang/Basic/DiagnosticSemaKinds.td

clang/lib/Sema/SemaDeclAttr.cpp

clang/test/CodeGenSYCL/Inputs/mock_properties.hpp

clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp

clang/include/clang/Basic/Attr.td

Fznamznon · 2022-10-04T11:00:04Z

clang/lib/Sema/SemaDeclAttr.cpp

+         diag::warn_sycl_old_and_new_kernel_attributes)
+        << ReqdWGSizeAttr;
+
+  if (const auto *ReqdSGSizeAttr = D->getAttr<IntelReqdSubGroupSizeAttr>())


Did I get it right that we issue a warning unconditionally if a kernel has any property and the three chosen attributes? Should we check which property is attached to a kernel?

That is right. The warning is issued if add_ir_attributes_function has values, regardless of the values. #6941 (comment) goes into more detail, but in summary we do it to avoid needing the frontend knowing about select values in this attribute while the warning is there to alert the user that they are using analogous features.

Ok, these three attributes are all that we can apply to SYCL kernel or is there more?

These are the only three added in this PR, but there will also be a device_has property in a follow-up. I suspect even more after that, but those are the ones defined by sycl_ext_oneapi_kernel_properties.

So, when there is a new property which conflicts with existing attribute, we have to update clang each time?

Sadly yes, though this should be rare after these, device_has, and any existing SYCL extensions with attributes, as we will be leaning towards adding compile-time properties extensions rather than extensions adding new attributes, avoiding any more conflicts.

clang/lib/Sema/SemaDeclAttr.cpp

clang/test/CodeGenSYCL/add_ir_attributes_function.cpp

Fznamznon · 2022-10-04T11:03:11Z

clang/include/clang/Basic/Attr.td

+        else if (const auto *IntegerVal =
+              dyn_cast<IntegerLiteral>(InitNoImpCastE))
+          C = static_cast<char>(IntegerVal->getValue().getZExtValue());
+        else


So, only character and integer are supported? Can part of property value be a double?

The design doc for compile-time properties allows floating point property values also, but I don't think we currently have any like that. If this PR is intending to implement the general support for compile-time properties, we should support FP types also:

https://github.com/intel/llvm/blob/sycl/sycl/doc/design/CompileTimeProperties.md#property-representation-in-c-attributes-and-in-ir

I think there's a bit of confusion here. This is only for parsing a string of characters defined as an initializer list. For this we should only need to be able to parse actual characters and integral values. Float, bool, etc are handled separately.

I see. But how we can do "support" (as you mentioned here #6941 (comment) ) of multi value property if we would like it to have a non-char/int as one of the values?

What we do in this case is that we convert a list of integers to a string with comma-separated values. For example 1, 2, 3 becomes "1,2,3".

Fznamznon · 2022-10-04T11:06:13Z

BTW, can we split this to separate patches that target only one component? I don't think that giant PRs with many different changes and reviewers is a good idea.

gmlueck · 2022-10-04T13:55:41Z

What is the plan for moving the extension specification out of "proposed" and into "experimental"? Do we want to make the APIs in this PR available soon, or do we plan to wait until the full extension is implemented? If we want to make them available before implementing the full extension, the extension specification should be split as described here.

I'd also like to update the extension specification to use the official template.

If it's easier, the extension spec can be split, moved, and updated in a separate PR.

steffenlarsen · 2022-10-04T19:33:53Z

BTW, can we split this to separate patches that target only one component? I don't think that giant PRs with many different changes and reviewers is a good idea.

Having it in one PR means we can do full E2E testing out-of-the-gate and since we now don't dismiss approvals when new changes are made I figured having the different components in the same would be okay as reviewers for the different components could focus on their select parts and potentially unsubscribe upon approval if the following messages are an annoyance. If you do not agree with that assessment I am okay with splitting it up in 3; sycl-post-link, cfe, and runtime headers. Keep in mind though that the testability of the last part would be severely limited until we have the other parts merged.

steffenlarsen · 2022-10-04T19:39:26Z

What is the plan for moving the extension specification out of "proposed" and into "experimental"? Do we want to make the APIs in this PR available soon, or do we plan to wait until the full extension is implemented? If we want to make them available before implementing the full extension, the extension specification should be split as described here.

I'd also like to update the extension specification to use the official template.

If it's easier, the extension spec can be split, moved, and updated in a separate PR.

The device_has part of the extension should follow soon, which is why this PR does not move the extension nor does it define the feature macro. The motivation for getting this in ASAP is to unblock development of other features leveraging the kernel properties infrastructure, independent of whether or not device_has is ready. If you think splitting it is preferred in the meantime, I am okay with that solution.

I'd also like to update the extension specification to use the official template.

I think it would make sense to do when we also move it.

gmlueck · 2022-10-04T19:41:57Z

If you think splitting it is preferred in the meantime, I am okay with that solution.

No need. I'm OK if we wait for device_has before considering the extension moved to "experimental".

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

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

sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp

Fznamznon · 2022-10-05T18:14:32Z

Having it in one PR means we can do full E2E testing out-of-the-gate and since we now don't dismiss approvals when new changes are made I figured having the different components in the same would be okay as reviewers for the different components could focus on their select parts and potentially unsubscribe upon approval if the following messages are an annoyance. If you do not agree with that assessment I am okay with splitting it up in 3; sycl-post-link, cfe, and runtime headers. Keep in mind though that the testability of the last part would be severely limited until we have the other parts merged.

What concerns me is that even though the approvals are retained there will be a lot of reviewers and therefore a ton of comments in the same PR. Because of that it will be hard to find previous conversations (for example if I want to see what are your replies on my comments). In order to perform full E2E testing you could open a separate draft PR which contains all changes and cherry pick commits that contain review comments from separate component-specific branches there. But If you still find it convenient, I'm not opposed to continue moving with this particular PR "as-is", but I would be happy to see my suggestion applied to future changes.

clang/include/clang/Basic/DiagnosticSemaKinds.td

clang/lib/Sema/SemaDeclAttr.cpp

clang/lib/Sema/SemaSYCL.cpp

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen · 2022-10-07T14:28:47Z

What concerns me is that even though the approvals are retained there will be a lot of reviewers and therefore a ton of comments in the same PR. Because of that it will be hard to find previous conversations (for example if I want to see what are your replies on my comments). In order to perform full E2E testing you could open a separate draft PR which contains all changes and cherry pick commits that contain review comments from separate component-specific branches there. But If you still find it convenient, I'm not opposed to continue moving with this particular PR "as-is", but I would be happy to see my suggestion applied to future changes.

That is a fair concern. I will split this into separate PRs ASAP and move this to draft. Any open discussions can either continue here or move. I will comment with links when it is done.

This commit makes the following changes to the Clang frontend in preparation for the initial implementation of the sycl_ext_oneapi_kernel_properties extension: * Fixed an issue preventing initializer lists of characters to be used as strings in add_ir_{attributes|annotations}_* attributes. * The add_ir_attributes_function will now be copied onto the generated kernel declaration from the recognized kernel_* functions. * Added a warning diagnostic for when a kernel has both an add_ir_attributes_function (with values) and one or more potentially conflicting SYCL attributes. This is split from intel#6941. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

This commit adds parsing of compile-time properties for kernels to sycl-post-link. This includes the following: * sycl-post-link is now able to generate metadata for recognized SYCL LLVM IR attributes on functions. These are "sycl-work-group-size", "sycl-work-group-size-hint", and "sycl-sub-group-size". * The previously mentioned new recognized LLVM IR attributes are translated into existing metadata the SPIR-V Translator translates into corresponding SPIR-V execution modes. If these metadata nodes already exist on a function (e.g. added through a SYCL 2020 attribute), the related property is ignored. This is split from intel#6941. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

This commit implements the following parts of the sycl_ext_oneapi_kernel_properties extension: * The work_group_size, work_group_size_hint, and sub_group_size properties. * The new overloads for single_task, parallel_for, and parallel_for_work_group in handler. * The new shortcuts for single_task and parallel_for in queue. * Support for merging property lists which is used when kernel functors have a get member for properties. The changes do not include any changes related to the device_has property and corresponding interfaces. This is split from intel#6941. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen · 2022-10-11T11:32:11Z

This has been split into three parts:

Clang frontend: [SYCL] CFE fixes for sycl_ext_oneapi_kernel_properties #7016
sycl-post-link: [SYCL] Add kernel property parsing to sycl-post-link #7017
SYCL headers: [SYCL] Implement subset of sycl_ext_oneapi_kernel_properties #7018 (Requires the other two PRs)

I will keep this open as draft for testing. Please feel free to move discussions to the other PRs or continue them here.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

This commit adds parsing of compile-time properties for kernels to sycl-post-link. This includes the following: * sycl-post-link is now able to generate metadata for recognized SYCL LLVM IR attributes on functions. These are "sycl-work-group-size", "sycl-work-group-size-hint", and "sycl-sub-group-size". * The previously mentioned new recognized LLVM IR attributes are translated into existing metadata the SPIR-V Translator translates into corresponding SPIR-V execution modes. If these metadata nodes already exist on a function (e.g. added through a SYCL 2020 attribute), the related property is ignored. This is split from #6941. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

This commit makes the following changes to the Clang frontend in preparation for the initial implementation of the sycl_ext_oneapi_kernel_properties extension: * Fixed an issue preventing initializer lists of characters to be used as strings in add_ir_{attributes|annotations}_* attributes. * The add_ir_attributes_function will now be copied onto the generated kernel declaration from the recognized kernel_* functions. * Added a warning diagnostic for when a kernel has both an add_ir_attributes_function (with values) and one or more potentially conflicting SYCL attributes. This is split from #6941. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

This commit implements the following parts of the sycl_ext_oneapi_kernel_properties extension: * The work_group_size, work_group_size_hint, and sub_group_size properties. * The new overloads for single_task, parallel_for, and parallel_for_work_group in handler. * The new shortcuts for single_task and parallel_for in queue. * Support for merging property lists which is used when kernel functors have a get member for properties. The changes do not include any changes related to the device_has property and corresponding interfaces. This is split from #6941. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen · 2022-10-18T18:54:11Z

All three parts have been merged. Thanks all!

Tests for intel/llvm#6941 Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Tests for intel#6941 Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen requested a review from Pennycook October 3, 2022 16:24

steffenlarsen mentioned this pull request Oct 3, 2022

[SYCL] Add tests for kernel properties intel/llvm-test-suite#1302

Merged

Remove triple from tests to avoid CUDA warning causing failures

996f47c

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Pennycook reviewed Oct 3, 2022

View reviewed changes

steffenlarsen added 2 commits October 3, 2022 14:46

Add warning diagnostic when properties and attributes are used

ba2e52c

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Fix formatting

ef20a0e

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen marked this pull request as ready for review October 3, 2022 21:52

steffenlarsen requested review from a team as code owners October 3, 2022 21:52

steffenlarsen requested a review from dm-vodopyanov October 3, 2022 21:52

Fix tests after formatting

7220e23

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Fznamznon reviewed Oct 4, 2022

View reviewed changes

steffenlarsen added 2 commits October 4, 2022 13:17

Respond to review comments

f2fd2af

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Use SExt instead of ZExt

cd215ce

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

AlexeySachkov reviewed Oct 5, 2022

View reviewed changes

elizabethandrews reviewed Oct 5, 2022

View reviewed changes

clang/include/clang/Basic/DiagnosticSemaKinds.td Outdated Show resolved Hide resolved

clang/lib/Sema/SemaDeclAttr.cpp Show resolved Hide resolved

clang/lib/Sema/SemaSYCL.cpp Show resolved Hide resolved

Respond to comments in sycl-post-link

adc707d

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

This was referenced Oct 11, 2022

[SYCL] CFE fixes for sycl_ext_oneapi_kernel_properties #7016

Merged

[SYCL] Add kernel property parsing to sycl-post-link #7017

Merged

[SYCL] Implement subset of sycl_ext_oneapi_kernel_properties #7018

Merged

steffenlarsen marked this pull request as draft October 11, 2022 11:32

Change variable name and use size_t analog from datalayout

7493f1f

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen added 7 commits October 13, 2022 04:47

Merge remote-tracking branch 'intel/sycl' into steffen/kernel_props

e8bc26d

Add hint to attr-add-ir-attributes-function-conflict

74856fd

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Add TODOs and a note about same-kernel-conflict

89f298e

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Fix formatting

0b40547

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Fix test after formatting

0a57078

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Change conflict warning

118ae2f

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Change diagnostic message further

10faff3

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

steffenlarsen closed this Oct 18, 2022

steffenlarsen added a commit to intel/llvm-test-suite that referenced this pull request Oct 19, 2022

[SYCL] Add tests for kernel properties (#1302)

74b271a

Tests for intel/llvm#6941 Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

aelovikov-intel pushed a commit to aelovikov-intel/llvm that referenced this pull request Mar 27, 2023

[SYCL] Add tests for kernel properties (intel/llvm-test-suite#1302)

9072987

Tests for intel#6941 Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[SYCL] Implement initial parts of sycl_ext_oneapi_kernel_properties #6941

[SYCL] Implement initial parts of sycl_ext_oneapi_kernel_properties #6941

steffenlarsen commented Oct 3, 2022 •

edited

Loading

Pennycook Oct 3, 2022

steffenlarsen Oct 3, 2022

steffenlarsen Oct 11, 2022

Pennycook Oct 11, 2022

gmlueck Oct 11, 2022

steffenlarsen Oct 13, 2022

asudarsa commented Oct 3, 2022

Fznamznon Oct 4, 2022

steffenlarsen Oct 4, 2022

Fznamznon Oct 5, 2022

steffenlarsen Oct 11, 2022

Fznamznon Oct 12, 2022

steffenlarsen Oct 12, 2022

Fznamznon Oct 4, 2022

gmlueck Oct 4, 2022

steffenlarsen Oct 4, 2022

Fznamznon Oct 5, 2022

steffenlarsen Oct 7, 2022

Fznamznon commented Oct 4, 2022

gmlueck commented Oct 4, 2022

steffenlarsen commented Oct 4, 2022

steffenlarsen commented Oct 4, 2022

gmlueck commented Oct 4, 2022

Fznamznon commented Oct 5, 2022

steffenlarsen commented Oct 7, 2022

steffenlarsen commented Oct 11, 2022

steffenlarsen commented Oct 18, 2022

		using NameT =
		typename detail::get_kernel_name_t<KernelName, KernelType>::name;

[SYCL] Implement initial parts of sycl_ext_oneapi_kernel_properties #6941

[SYCL] Implement initial parts of sycl_ext_oneapi_kernel_properties #6941

Conversation

steffenlarsen commented Oct 3, 2022 • 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

asudarsa commented Oct 3, 2022

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

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

Fznamznon commented Oct 4, 2022

gmlueck commented Oct 4, 2022

steffenlarsen commented Oct 4, 2022

steffenlarsen commented Oct 4, 2022

gmlueck commented Oct 4, 2022

Fznamznon commented Oct 5, 2022

steffenlarsen commented Oct 7, 2022

steffenlarsen commented Oct 11, 2022

steffenlarsen commented Oct 18, 2022

steffenlarsen commented Oct 3, 2022 •

edited

Loading