[libspirv] Use the remangler to mangle half types #18083
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Commit 316418d modified the AMDGCN sub-group shuffle builtins. In doing so it broke them in a couple of ways. The first is a typo that hid the SubgroupShuffleXorINTEL builtins for sub 32-bit types but ommitting the 'INTEL' part of the builtin name. The second was an issue of mangling for the 'half' data type, which accidentally stopped the SYCL host being able to find the right symbols. Since libspirv is defined in OpenCL C, the plain 'half' data type is mangled differently to the sycl::half type used by SYCL: 'Dh' vs 'DF16_'. We've long worked around this using manual mangling of the builtins, which has never been very nice and has often caused bugs that are awkward to track down. This PR takes an alternative approach and instead uses a custom typedef to the _Float16 type, and defines the sub-group shuffle builtins using that new type. Since _Float16 is the 'half' type used everywhere except OpenCL, we arrive at the correct mangling. If this approach proves successful we might be able to expand on it and remove much of the manual mangling elsewhere in libspirv.
frasercrmck
had a problem deploying
to
WindowsCILock
GitHub Actions
Error
ldrumm
reviewed
Apr 17, 2025
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
ldrumm
approved these changes
Apr 17, 2025
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
changed the title
frasercrmck
changed the title
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
frasercrmck
temporarily deployed
to
WindowsCILock
GitHub Actions
Inactive
Contributor
Author
|
@intel/llvm-gatekeepers This is ready to merge, thanks. |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This commit revisits the way we achieve the desired host-side mangling of
halftypes asDF16_(from the_Float16type) coming from the original mangling ofDhtype that OpenCL produces. We were previously manually achieving this by writing over a thousand wrapper functions from_Float16types tohalftypes.The remangler can just as easily do this for us, while reducing the total source code line count by over 8000.
This work was originally spurred on by commit 316418d, which modified the AMDGCN sub-group shuffle builtins. In doing so it broke them in a couple of ways.
The first is a typo that hid the SubgroupShuffleXorINTEL builtins for sub 32-bit types but ommitting the 'INTEL' part of the builtin name. That was easily fixed by adjusting the builtin name.
The second was an issue of providing builtins using the OpenCL 'half' data type, without the
_Float16->halfwrappers which we relied on. This accidentally stopped the SYCL host being able to find the right symbols.