Description
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.