cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

stefan_w
Journeyman III

CL_FILTER_LINEAR not working

Samplers with CL_FILTER_LINEAR seem to use nearest neighbor filtering instead of bilinear interpolation. Or can anybody confirm that bilinear filtering with OpenCL is working on ATI cards?

0 Likes
6 Replies
genaganna
Journeyman III

Originally posted by: stefan_w Samplers with CL_FILTER_LINEAR seem to use nearest neighbor filtering instead of bilinear interpolation. Or can anybody confirm that bilinear filtering with OpenCL is working on ATI cards?

 

Yes it works.  Please see SobelFilterImage for reference.  Please send us your code to reproduce the issue.

0 Likes

The SobelFilterImage example uses only integer coordinates. Therefore there is no difference between nearest neighbor and linear filtering.

try the following modified version of SimpleImage_Kernels.cl and test if there is a difference between CLK_FILTER_LINEAR and CLK_FILTER_NEAREST.

 

__constant sampler_t imageSampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR; /* Copy input 2D image to output 2D image */ __kernel void image2dCopy(__read_only image2d_t input, __write_only image2d_t output) { float2 coord = (float2)((float)get_global_id(0)*0.01f, get_global_id(1)); int2 coordout = (int2)(get_global_id(0), get_global_id(1)); uint4 temp = read_imageui(input, imageSampler, coord); write_imageui(output, coordout, temp); } /* Copy input 3D image to 2D image */ __kernel void image3dCopy(__read_only image3d_t input, __write_only image2d_t output) { int2 coord = (int2)(get_global_id(0), get_global_id(1)); /* Read first slice into lower half */ uint4 temp0 = read_imageui(input, imageSampler, (int4)(coord, 0, 0)); /* Read second slice into upper half */ uint4 temp1 = read_imageui(input, imageSampler, (int4)((int2)(get_global_id(0), get_global_id(1) - get_global_size(1)/2), 1, 0)); write_imageui(output, coord, temp0 + temp1); }

0 Likes

Linear filtering works with read_imagef function as mentioned in the spec. I can clearly see the visual difference between nearest and linear filtering modes in my volume rendering app.

 

0 Likes

Thank you. Bilinear filtering is no working for 8bit and 16 bit normalized integer Textures. However with textures of type CL_FLOAT bilinear interpolation does not work for me. The spec says that bilinear interpolation on CL_FLOAT textures is not working in the embedded profile, so I think it should work in full profile mode.

0 Likes

stefan_w,

Can you please provide test case.It is working on our system.

You can send your test case to streamdeveloper@amd.com if the code is proprietary.

Thanks

0 Likes

stefan_w,

As per spec:
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.

So CLK_FILTER_LINEAR is only expected to work for read_imagef and not for
read_imageui.

Is your confusion resolved? If no can you post the host code also.

Thanks

0 Likes