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

Refactor vector_swizzles test #825

Conversation

AlexeySachkov
Copy link
Contributor

This is a partial refactoring of the test, which aims to:

  • reduce kernel size, which allows to speed up JIT compilation and therefore test execution
  • improve error reporting from the test, by returning more than just a single bool value for all of the checks performed by kernels

Main changes:

  • introduced new template for swizzle test kernels
  • results of "per-element" .swizzle<s0, s1, ..., sN>() operations are now reported back to host where their verification happens instead of doing so on device side

This is a partial refactoring of the test, which aims to:
- reduce kernel size, which allows to speed up JIT compilation and
  therefore test execution
- improve error reporting from the test, by returning more than just a
  single boolean value for all of the checks performed by kernels

Main changes:
- introduced new template for swizzle test
- "basic" `.swizzle<>()` checks now report results back to host where
  their verification happens instead of doing so on device side
@AlexeySachkov AlexeySachkov requested a review from a team as a code owner October 31, 2023 20:27
if (!check_equal_type_bool<sycl::vec<${type}, ${size}>>(inOrderSwizzleFunctionVec)) {
resAcc[0] = false;
}
if (!check_vector_size<${type}, ${size}>(inOrderSwizzleFunctionVec)) {
resAcc[0] = false;
}
if (!check_vector_values<${type}, ${size}>(inOrderSwizzleFunctionVec, in_order_vals)) {
Copy link
Contributor Author

Choose a reason for hiding this comment

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

Note for reviewers: there are two templates for kernel code:

  • one for vectors with 4 or less elements, which is used to exercise all simple swizzles like .xyzw; modified swizzle_elem_template is its part
  • another one for bigger vectors, where only per-element swizzles are tested; swizzle_full_test_template is used for it

Kernels for both categories of vectors do the same checks to ensure that .swizzle method returns elements in the right order, they are just described by two different code templates.

@AlexeySachkov
Copy link
Contributor Author

Ideally, we should get a better rewrite of the test to be able to emit separate error message for (almost) every check performed by the test, so SYCL implementors can quickly understand which exact operation doesn't work properly. Plus, we do not test swizzles functionality on host where it should also work, according to the SYCL 2020 specification.

However, I'm not able to provide such a rewrite at the moment. Nevertheless, I would appreciate if this partial refactoring is accepted, as it greatly helps to reduce to reduce amount of code emitted for device. I don't have pure numbers, because I have some local changes applied to intel/llvm implementation which also help this particular test, but I would estimate that this patch reduces device code size by 30-40% when CTS is compiled by intel/llvm implementation.

@bader bader added the help wanted Extra attention is needed label Oct 31, 2023
Copy link
Member

@keryell keryell left a comment

Choose a reason for hiding this comment

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

Thanks.
You can probably add your idea for further refactoring as a TODO comment somewhere so it is not lost.

@AlexeySachkov
Copy link
Contributor Author

Thanks. You can probably add your idea for further refactoring as a TODO comment somewhere so it is not lost.

Recorded key points in 35fbe65

Copy link
Member

@keryell keryell left a comment

Choose a reason for hiding this comment

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

Thanks for the on-coming generations of developers. :-)

@bader bader merged commit 75b288a into KhronosGroup:SYCL-2020 Nov 9, 2023
7 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
help wanted Extra attention is needed
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants