Community
cancel
Showing results for 
Search instead for 
Did you mean: 
Juan_Ignacio_C_
Beginner
65 Views

CLK_FILTER_LINEAR not working in version 15.1 under Linux

Reading an image using CLK_FILTER_LINEAR interpolation mode produces zeros instead of the correct image values.

Driver: OpenCL Runtime 15.1

Version: 5.0.0.57

OS: Debian Jessie

Language: Python (PyOpenCL)

The following kernel is a minimum test case. It is supposed to copy the input into the output. It works fine under the Beignet GPU driver, but under Intel's CPU driver the output image consists of all zeros. When CLK_FILTER_LINEAR is replaced by CLK_FILTER_NEAREST it works.

The problem persists when using normalized coordinates.

constant sampler_t S =  CLK_NORMALIZED_COORDS_FALSE
                        | CLK_ADDRESS_NONE | CLK_FILTER_LINEAR;

kernel void tube(
        read_only image2d_t img_in,
        write_only image2d_t img_out)
{
    int2 xy = {get_global_id(0), get_global_id(1)};
    uint4 rgb_;

    rgb_ = read_imageui(img_in, S, convert_float2(xy) + 0.5f);
    write_imageui(img_out, xy, rgb_);
}

I attach a small test script.

0 Kudos
6 Replies
Robert_I_Intel
Employee
65 Views

Hi Juan Ignacio,

 Sorry for the delay: I crashed my Linux box after I installed a bunch of dependencies required by your python script - will take me some time to reinstall everything and try again.

Juan_Ignacio_C_
Beginner
65 Views

Could you reproduce the issue? I can make a C example if you are having problems with python.

Robert_I_Intel
Employee
65 Views

I wasn't able to reproduce this with Python, so yes, a C example would help.

I talked with the CPU runtime lead and he acknowledged that there are problems with sampler implementation on the CPU, but was not sure whether it was this specific issue or something else, so having a C example would allow me to file a bug and track its fix.

Thanks!

Juan_Ignacio_C_
Beginner
65 Views

Here is a C example. It saves the input and output image to PGM format. Additionally, it computes the difference between both images.

Robert_I_Intel
Employee
65 Views

Hi Juan Ignacio,

Yes, I can reproduce it on a CPU device. Thank you for providing such a beautiful reproducer! I will notify the CPU runtime team.

Thanks!

Robert_I_Intel
Employee
65 Views

Hi Juan Ignacio,

The member of the CPU team pointed me to the following:

From https://www.khronos.org/registry/cl/specs/opencl-1.2.pdf p.298

The read_image{i|ui} calls support a nearest filter only. The filter_mode specified in sampler must be set to CLK_FILTER_NEAREST; otherwise the values returned are undefined.

CPU is sensitive to such kind of bugs because it is undefined behavior. on CPU it may lead to crash due to Access violation since OS takes care of that. On GPU Access violations, sampler misusage are more tolerant. App is not terminated.