6 Replies Latest reply on Dec 29, 2010 5:57 AM by himanshu.gautam

    CL_FILTER_LINEAR not working

    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?

        • CL_FILTER_LINEAR not working
          genaganna

           

          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.

            • CL_FILTER_LINEAR not working
              stefan_w

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