Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[SYCL] Add experimental implementation of sycl_ext_intel_grf_size #9882

Merged
merged 11 commits into from
Jun 28, 2023
Merged

[SYCL] Add experimental implementation of sycl_ext_intel_grf_size #9882

merged 11 commits into from
Jun 28, 2023

Conversation

sarnex
Copy link
Contributor

@sarnex sarnex commented Jun 14, 2023 •
edited
Loading

This change implements sycl_ext_intel_grf_size, and in particular:

  1. Moves the sycl_ext_intel_grf_size spec document from the proposed folder to the experimental folder, and updates the implementation status in the document to match.
  2. Adds two new kernel properties sycl::ext::intel::experimental::grf_size and sycl::ext::intel::experimental::grf_size_automatic, as per the spec. grf_size adds the sycl-grf-size metadata with a value of the template parameter (128 or 256). grf_size_automatic adds the sycl-grf-size metadata with a value of 0.
  3. Marks the sycl::detail::register_alloc_mode property as deprecated, and it still works as before.
  4. Updates CompileTimePropertiesPass.cpp to map the sycl-grf-size metadata added by the front-end to the RegisterAllocMode metadata which llvm-spirv looks for. This RegisterAllocMode metadata is how AOT works.
  5. Updates sycl-post-link to split by the sycl-grf-size metadata, have a sycl-grf-size binary property, and do a error check to make sure the deprecated sycl::detail::register_alloc_mode and grf_size/grf_size_automatic are not set at the same time on the same kernel.
  6. Updates program_manager to deal with the new image property and pass the right flags
  7. Updates exists tests to also test the new properties.

@sarnex sarnex temporarily deployed to aws June 14, 2023 17:52 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 14, 2023 19:10 — with GitHub Actions Inactive
@sarnex sarnex marked this pull request as ready for review June 14, 2023 20:57
@sarnex sarnex requested review from a team as code owners June 14, 2023 20:57
@sarnex sarnex temporarily deployed to aws June 15, 2023 14:12 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 15, 2023 14:51 — with GitHub Actions Inactive
v-klochkov
v-klochkov reviewed Jun 15, 2023
View reviewed changes
llvm/lib/SYCLLowerIR/CompileTimePropertiesPass.cpp Outdated Show resolved Hide resolved
llvm/test/SYCLLowerIR/CompileTimePropertiesPass/kernel-attributes/grf-size.ll Show resolved Hide resolved
sycl/test-e2e/ESIMD/large-grf.cpp Show resolved Hide resolved
sycl/test-e2e/DeviceCodeSplit/grf.cpp Outdated Show resolved Hide resolved
sycl/include/sycl/detail/kernel_properties.hpp Outdated Show resolved Hide resolved
llvm/tools/sycl-post-link/sycl-post-link.cpp Outdated Show resolved Hide resolved
llvm/test/tools/sycl-post-link/sycl-grf-size.ll Outdated Show resolved Hide resolved
@sarnex sarnex temporarily deployed to aws June 16, 2023 19:17 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 16, 2023 20:02 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 16, 2023 20:43 — with GitHub Actions Inactive
asudarsa
asudarsa reviewed Jun 19, 2023
View reviewed changes
Copy link
Contributor

@asudarsa asudarsa left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I have a question about long-term implementation support for this. dpc++ tools changes look good. Please address comment.

Thanks

@sarnex sarnex temporarily deployed to aws June 20, 2023 15:36 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 20, 2023 21:42 — with GitHub Actions Inactive
steffenlarsen
steffenlarsen reviewed Jun 21, 2023
View reviewed changes
sycl/include/sycl/detail/kernel_properties.hpp
inline constexpr register_alloc_mode_key::value_t<Mode> register_alloc_mode
__SYCL_DEPRECATED("register_alloc_mode is deprecated, "
"use sycl::ext::intel::experimental::grf_size or "
"sycl::ext::intel::experimental::grf_size_automatic");
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

@stdale-intel - Do we want to deprecate this? It has not been in a release yet, to my knowledge.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

From my POV at least we shouldn't remove it right now, otherwise all of the clients who just ported to using it would have to make yet another change to their code, annoying them even further.

sycl/include/sycl/ext/oneapi/properties/property.hpp Show resolved Hide resolved
@sarnex sarnex temporarily deployed to aws June 21, 2023 20:11 — with GitHub Actions Inactive
@sarnex
Copy link
Contributor Author

sarnex commented Jun 21, 2023

Sorry for the force-push, my original HEAD had CI failing even with no changes.

@sarnex sarnex temporarily deployed to aws June 21, 2023 21:03 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 21, 2023 21:51 — with GitHub Actions Inactive
@sarnex
Copy link
Contributor Author

sarnex commented Jun 26, 2023

@intel/dpcpp-tools-reviewers @AlexeySachkov ping on this one? thanks!

@sarnex
Copy link
Contributor Author

sarnex commented Jun 27, 2023 •
edited
Loading

@steffenlarsen do you mind formally approving this pr? my impression is you are alright with it (hopefully that's not a misread :) ), we technically have runtime approval from slava, but i don't think he is working closely in the runtime area at the moment.

steffenlarsen
steffenlarsen approved these changes Jun 27, 2023
View reviewed changes
Copy link
Contributor

@steffenlarsen steffenlarsen left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM!

@sarnex sarnex temporarily deployed to aws June 27, 2023 15:03 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 27, 2023 15:44 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 27, 2023 17:17 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 27, 2023 18:12 — with GitHub Actions Inactive
sarnex and others added 9 commits June 27, 2023 12:32
@sarnex
[SYCL] Add experimental implementation of sycl_ext_intel_grf_size
f8991b1 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
typo
eaa6b3d 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
Make sentence bold
c9e55b3
@sarnex
Address feedback
4858edf 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
improve lit test
6c128ef 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
Add TODO comment
f362bc9 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
Add conflicting property checking infrastructure
55e3001 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
Add compile time assert for invalid GRF sizes
139fe05 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex
Update LIT test to pass on HEAD
755419a 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex sarnex temporarily deployed to aws June 27, 2023 19:53 — with GitHub Actions Inactive
@sarnex
random commit to try to get CI to start
edc9b6e 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex sarnex temporarily deployed to aws June 27, 2023 20:22 — with GitHub Actions Inactive
@sarnex
fix LIT test again
010955c 
Signed-off-by: Sarnie, Nick <nick.sarnie@intel.com>
@sarnex sarnex temporarily deployed to aws June 27, 2023 21:07 — with GitHub Actions Inactive
@sarnex sarnex temporarily deployed to aws June 27, 2023 21:47 — with GitHub Actions Inactive
@sarnex
Copy link
Contributor Author

sarnex commented Jun 28, 2023

@intel/llvm-gatekeepers This one is ready for merge, thanks!

@steffenlarsen steffenlarsen merged commit 370aa2a into intel:sycl Jun 28, 2023
@Chenyang-L Chenyang-L mentioned this pull request Jul 11, 2023
Closed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@AlexeySachkov AlexeySachkov AlexeySachkov approved these changes

@asudarsa asudarsa asudarsa left review comments

@gmlueck gmlueck gmlueck approved these changes

@v-klochkov v-klochkov v-klochkov approved these changes

@steffenlarsen steffenlarsen steffenlarsen approved these changes

Assignees
No one assigned
Labels
None yet
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
6 participants
@sarnex @gmlueck @AlexeySachkov @v-klochkov @asudarsa @steffenlarsen