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.

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

Improve check for compatibility of vector size and subgroup size in SYCL #4579

Merged

Conversation

masterleinad
Copy link
Contributor

@masterleinad masterleinad commented Dec 2, 2021

Blocked by intel/llvm#5071.

@masterleinad masterleinad force-pushed the improve_check_vector_size_subgroup_size branch from b97b08e to cff2a34 Compare December 3, 2021 16:40
@masterleinad
Copy link
Contributor Author

This should work now. @brian-kelley Can you give this a try in view of #4573?

@masterleinad masterleinad marked this pull request as ready for review December 3, 2021 16:42
@masterleinad
Copy link
Contributor Author

Retest this please.

@brian-kelley
Copy link
Contributor

@masterleinad Yes, this does catch the vector size being too big.

@masterleinad
Copy link
Contributor Author

Retest this please.

1 similar comment
@masterleinad
Copy link
Contributor Author

Retest this please.

@brian-kelley
Copy link
Contributor

brian-kelley commented Dec 8, 2021

@masterleinad Is "subgroup size" equivalent to team size for Kokkos? Is this PR expected to catch issues with team size being too big? Coincidentally, the same kernel that had the local atomics issue now hangs because policy.team_size_recommended() returns 1024 which I guess is too big. Reducing it to 256 makes it work fine. If that is unrelated to this PR, does that deserve its own issue? The policy uses vector_length 1.

@masterleinad
Copy link
Contributor Author

@masterleinad Is "subgroup size" equivalent to team size for Kokkos? Is this PR expected to catch issues with team size being too big? Coincidentally, the same kernel that had the local atomics issue now hangs because policy.team_size_recommended() returns 1024 which I guess is too big. Reducing it to 256 makes it work fine. If that is unrelated to this PR, does that deserve its own issue? The policy uses vector_length 1.

"group size" is equivalent to team size and subgroup size is related to vector size. I haven't seen a kernel hanging because of the team size/subgroup size being too large; normally there is some kind of error coming from the runtime. Anyway, please open a separate issue!

@masterleinad masterleinad force-pushed the improve_check_vector_size_subgroup_size branch from 3a036b5 to 7d0d595 Compare December 8, 2021 17:44
@masterleinad
Copy link
Contributor Author

This seems to finally work with storing the kernel obtained from the runtime in a static variable and also doing another warm-up round in the test making sure that obtaining the kernel is not counted when testing for the kernel execution to be asynchronous.

@masterleinad masterleinad force-pushed the improve_check_vector_size_subgroup_size branch from 0861ad1 to 0d64ed3 Compare December 9, 2021 18:05
Comment on lines +447 to +448
kernel
.get_info<sycl::info::kernel_device_specific::max_sub_group_size>(
Copy link
Member

Choose a reason for hiding this comment

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

Format looks weird

@dalg24 dalg24 merged commit 6212cc0 into kokkos:develop Dec 14, 2021
@masterleinad masterleinad deleted the improve_check_vector_size_subgroup_size branch December 14, 2021 20:02
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

4 participants