-
Notifications
You must be signed in to change notification settings - Fork 733
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
[SYCL] Fix group algorithms for non-uniform groups, marray and vec #14364
[SYCL] Fix group algorithms for non-uniform groups, marray and vec #14364
Conversation
The implementation currently defines the GroupNonUniform SPIR-V builtins with the same extension suffixes as their Group counterparts. However, some of the Group counterparts are KHR extension functions, while the GroupNonUniform operations are core SPIR-V operations. As such, the SYCL-CTS tests fail to compile the non-uniform group extension tests in full-conformance mode. This commit fixes this issue. Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Can we add at least a compile-only in-tree test to have better in-tree coverage?
While writing compile-only testing, more issues with group algorithms have appeared. I will convert this to draft while I address these. |
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
A test for shuffles with marray has been added. The rest should be covered by the compile-time test, so this is ready for review. @bso-intel - Could you please have another look? |
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
LGTM
…ntel#14364) This commit makes the following fixes to group algorithms: * Some GroupNonUniform SPIR-V builtins were incorrectly named as they were bundled together with their KHR Group SPIR-V equivalents. These have been renamed to map correctly to the right SPIR-V operations. * `sycl::marray` is now considered when checking for arithmetic types, making it usable in group broadcast operations and the functions that use them. * The representation of `bool` in `sycl::vec` has been changed to unsigned to match the representation picked by `ConvertToOpenCLType`. * The representation of `char` in `sycl::vec` has been changed to to match the representation picked by `ConvertToOpenCLType` to avoid cases where signedness would cause type mismatches. --------- Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
This commit makes the following fixes to group algorithms:
sycl::marray
is now considered when checking for arithmetic types, making it usable in group broadcast operations and the functions that use them.bool
insycl::vec
has been changed to unsigned to match the representation picked byConvertToOpenCLType
.char
insycl::vec
has been changed to to match the representation picked byConvertToOpenCLType
to avoid cases where signedness would cause type mismatches.