Kernel not working anymore with latest driver.

[LLess]

I would like to check if anybody else is seeing things like that.

I have a simple kernel doing a simple texture copy (npower2 to p2) conversion that was working fine with OpenCL 1.1 and with OpenCL 1.2 I am getting a black image.

Other kernels seems to be working fine so that points at a bug somewhere but I have strictly no idea where.

//Texturing into a power of 2 target
float4 RunPass ( __read_only image2d_t input_texture, int2 pos, __constant ImageFilterParams *params)
{
    int2 input_dim =get_image_dim(input_texture);
    pos.x = pos.x - params->outputTextureCenterX + (input_dim.x/2);
    pos.y = pos.y - params->outputTextureCenterY + (input_dim.y/2);
    if ( pos.x < 0 || pos.x >= input_dim.x || pos.y < 0 || pos.y >input_dim.y )
        return (float4)(0.0f);   
    return read_imagef(input_texture,int_nearest_sampler2, pos );
}

__kernel void ImageFilterKernel(__read_only image2d_t texture, __write_only image2d_t output, __constant ImageFilterParams *params)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int2 image_dim =get_image_dim(output);
    if  (( y < image_dim.y ) && (x < image_dim.x ))
    {
        int2 position = (int2)(x,y);
        float4 color = RunPass(texture,position,params);
        write_imagef(output,position,color);
    }
}

No failure reported by the functions that I am calling.

Laurent

[LLess]

//Texturing into a power of 2 target
float4 RunPass ( __read_only image2d_t input_texture, int2 pos, int2 outputCenter)
{
    int2 input_dim =get_image_dim(input_texture);
    int2 p;
    p.x = pos.x - outputCenter.x + (input_dim.x/2);
    p.y = pos.y - outputCenter.y + (input_dim.y/2);
    if ( p.x < 0 || p.x >= input_dim.x || p.y < 0 || p.y >input_dim.y )
        return (float4)(0.0f);
   
    return read_imagef(input_texture,int_nearest_sampler2, p );
}

__kernel void ImageFilterKernel(__read_only image2d_t texture, __write_only image2d_t output)
{
    int x = get_global_id(0);
    int y = get_global_id(1);
    int2 image_dim =get_image_dim(output);
    int2 image_center_dim = (int2)(image_dim.x/2,image_dim.y/2);
    if  (( y < image_dim.y ) && (x < image_dim.x ))
    {
        int2 position = (int2)(x,y);
        float4 color = RunPass(texture,position,image_center_dim);
        write_imagef(output,position,color);
    }
}

But this works...

Go figure.

[Raghupathi_M_Intel]

Laurent,

The kernel outputs incorrect image if you pass the ImageFilterParams* as a parameter but works ok if you compute the center in your kernel and pass to RunPass? Can you share your host code too?

Thanks,
Raghu

[LLess]

I can probably do that.

I will email you the code of that class and the other one needed but it is will be difficult to send you something that compiles since it would need almost the full project.

Laurent

[LLess]

I can now confirm that reverting back to the driver supporting only OpenCL 1.1 on the HD4000 fixed all my troubles.

Laurent.