cancel
Showing results for 
Search instead for 
Did you mean: 

Archives Discussions

s58000
Adept II

filter mode CLK_FILTER_LINEAR

CLK_FILTER_LINEAR seems not to work on HD 6950

I've tried to to use the CLK_FILTER_LINEAR filter mode with read_imagef()  on a 2d-Image without success, i get the same results when using CLK_FILTER_NEAREST.

Tested with the 2.4 SDK and 11.3 driver on the Radeon HD 6950.

Can anyone confirm??

0 Likes
12 Replies
himanshu_gautam
Grandmaster

How do you conclude that bahaviour. Can you please post a test case.

0 Likes

I've created the image with "clCreateImage2D()" using CL_FLOAT and CL_RGBA format.
I can confirm that following sampler combinations with "read_imagef()" still return CLK_FILTER_NEAREST results with both int2 and float2 coordinates:

__constant sampler_t samp = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_NEAREST | CLK_ADDRESS_CLAMP;
__constant sampler_t sampf = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP;

__kernel void advect(__write_only image2d_t    destImage,        
                     __read_only image2d_t    srcImage,        
                     int2                    size            )
{
    int2 pos = (int2)(get_global_id(0), get_global_id(1));
    if(pos.x < size.x && pos.y < size.y)
    {
        float2 posf = (float2)(pos.x + 0.5f, pos.y + 0.5f);
        float2 new_posf = posf - read_imagef(srcImage, samp, pos).xy;
        float4 val = read_imagef(srcImage, sampf, new_posf);
        write_imagef(destImage, pos, val);
    }
}

Same thing for 3D images.


and I can confirm that i'm getting the right results when using the same kernels on several NVIDIA GPUs.

0 Likes

Am I the only one having this issue ???

0 Likes

anyone ???????

0 Likes

@MicahVillmow

No problem ... thank you...

0 Likes

any news regarding this issue ??

0 Likes

sorry for the late reply.

Can you please post the host code also. Although i am not very sure but i think in some cases the results obtained from the two filters can be same.

It would be very easy to verify if it is a bug after you send the complete code.

Thanks

0 Likes

@ himanshu

unfortunately i can't post the host code here because it is a part of a big project and i'm using QtOpenCL wrapper. All the configuration i used for the creation of the 2D Image are posted above. I've tried all this combinations already:

__constant sampler_t samp_n = CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP;
__constant sampler_t sampf = CLK_NORMALIZED_COORDS_FALSE | CLK_FILTER_LINEAR | CLK_ADDRESS_CLAMP;

1) i've used normalized coordinates with float2 in the [0 - 1] range

2) unnormalized coordinates with float2 or int2 [0 - imageWidth/imageHeight]

3) tried also using CLK_ADDRESS_CLAMP / CLK_ADDRESS_CLAMP_TO_EDGE / CLK_ADDRESS_NONE

the result is always the same as the CLK_FILTER_NEARES. Same thing for 3D Images !!!

AND ALL this combinations are doing fine on 4 different NVIDIA GPUs.

I hope this will be solved in the next release. Thank you.

0 Likes

I'm facing with same problem. (using HD5750)

Seeing sources in AMD APP samples, 

texture fetching is executed by

float4 read_imagef (image2d_t image, sampler_t sampler, int2 coord)

but not

float4 read_imagef (image2d_t image, sampler_t sampler, float2 coord).

Does it mean that CLK_FILTER_LINEAR  isn't supprted in current driver ?

Fetching with CLK_FILTER_LINEAR works properly in NVIDIA's GPU.

0 Likes

any news about this ???

0 Likes

Fixed with Catalyst 11.7 !!!

thank you

0 Likes

s5800,
Sorry that I missed this, we've been really busy recently. I'm having someone look into it.
0 Likes