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?
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.
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); }
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.
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.
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
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