Read on RGBA unrom_int8 Image Always 0 with uint4

[peiyigu-intel]

I compile the following code with 2024.0 oneAPI base tool kit.
.\compiler\2024.0\bin\icpx.exe -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend "-device dg2" scaling.cpp -o scaling.exe
It seems the read pixel is always 0 when using uint4. Only can read data when using half4 which is float type

#include <CL/sycl.hpp>
#include <iostream>


using namespace cl::sycl;

class Scaling
{
public:
    Scaling(accessor<uint4, 2, access::mode::read, access::target::image>  inputImg,
        accessor<uint4, 2, access::mode::write, access::target::image> outputImg,
        sampler sampler) :
        m_inputImage(inputImg), m_outputImage(outputImg), m_sampler(sampler)
    {}

    void operator()(sycl::item<2> idx) const
    {
        float2 coord((float)idx[0] / (float)m_outputImage.get_range()[0] , (float)idx[1] / (float)m_outputImage.get_range()[1]);
        auto pixel = m_inputImage.read(coord, m_sampler);
        uint4 pp = {};
        pp.x() = static_cast<uint8_t>(125);
        pp.y() = static_cast<uint8_t>(122);
        pp.z() = static_cast<uint8_t>(121);
        pp.w() = static_cast<uint8_t>(199);
        m_outputImage.write(int2(idx[0], idx[1]), pp);
        //m_debugInfo[0] = pixel.x();
    }

    accessor<uint4, 2, access::mode::read, access::target::image> m_inputImage;
    accessor<uint4, 2, access::mode::write, access::target::image> m_outputImage;
    //unsampled_image_accessor<uint4, 2, access::mode::write> m_outputImage;
    sampler m_sampler;

    //accessor<int, 1, access::mode::write> m_debugInfo;
};

int main() {
    std::cout << "start process" << std::endl;

    uint4 tmpIn[100];
    uint4 tmpOut[400];
    for (int i = 0; i < 100; i++) {
        tmpIn[i] = uint4(0,0,0,0);
        tmpOut[i] = uint4(0, 0, 0, 0);
    }


    try {
        image<2> srcImage(tmpIn, image_channel_order::rgba, image_channel_type::unorm_int8, range<2>(10, 10));
        image<2> destImage(tmpOut, image_channel_order::rgba, image_channel_type::unorm_int8, range<2>(20, 20));

        //sycl::unsampled_image<2> destImage(tmpOut, sycl::image_format::r8g8b8a8_unorm, sycl::range<2>(10, 10));

        //buffer<int, 1> buffer(debug_info.data(), debug_info.size());

        std::cout << "image load done" << std::endl;

        queue q;
        std::cout << "Running on "
            << q.get_device().get_info<sycl::info::device::name>()
            << "\n";

        q.submit([&](handler& cgh) {
            accessor<uint4, 2, access::mode::read, access::target::image> image_acc(srcImage, cgh);
            auto output_acc = destImage.get_access<uint4, access::mode::write>(cgh);

            sampler smplr(coordinate_normalization_mode::unnormalized,
                addressing_mode::clamp,
                filtering_mode::linear);

            std::cout << "accessor ready" << std::endl;

            cgh.parallel_for<Scaling>(range<2>(20, 20), Scaling(image_acc, output_acc, smplr));

            }).wait();
    }
    catch (exception e)
    {

    }

    std::cout << "gpu process done" << std::endl;

    std::cout << "output: " << tmpOut[0].r() << " " << tmpOut[0].b() << " " << tmpOut[0].g() << " " << tmpOut[0].a() << std::endl;
    float a = tmpOut[10].a();
    std::cout << "output a :" << int(a) << std::endl;

    std::cout << "all done" << std::endl;

    return 0;
}