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

[SYCL][L0] Fix __spirv_ImageWrite impl so that read_write_user_type.cpp passes #17807

Open
JackAKirk opened this issue Apr 2, 2025 · 0 comments
Labels
bug Something isn't working SPIR-V Issues related to SPIRV-LLVM-Translator

Comments

@JackAKirk
Copy link
Contributor

JackAKirk commented Apr 2, 2025

Describe the bug

test-e2e/bindless_images/user_types/read_write_user_type.cpp
has test failures only on l0 and only for the char/short/half user defined vector types.
This occurs on both arc-a770 and BMG.

Upon further investigation, by bypassing write_image_array and writing myPixel directly to host accessible usm, it turns out that the fetch_image and sample_image APIs are working correctly.
The problem appears to be isolated to write_image_array, and by checking sycl::vec using char/short/half types I can also isolate it to user_defined types as covered by this test.
Adjusting the implementation of detail::convert_color shows that the failures/passes do not depend upon the particular implementation of the bitcasting implementation from the user types. Therefore
my conclusion is that the problem is likely due to a broken implementation of __spirv_ImageWrite (called via __invoke__ImageWrite) in the L0 backend.

@JackAKirk JackAKirk added bug Something isn't working SPIR-V Issues related to SPIRV-LLVM-Translator labels Apr 2, 2025
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
bug Something isn't working SPIR-V Issues related to SPIRV-LLVM-Translator
Projects
None yet
Development

No branches or pull requests

1 participant