Skip to content

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

Open
@JackAKirk

Description

@JackAKirk

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.

Metadata

Metadata

Assignees

No one assigned

    Labels

    SPIR-VIssues related to SPIRV-LLVM-TranslatorbugSomething isn't workinglevel-zeroIssues related to the Level Zero backendsycl-bindless-imagesSYCL Bindless Images

    Type

    No type

    Projects

    No projects

    Milestone

    No milestone

    Relationships

    None yet

    Development

    No branches or pull requests

    Issue actions