OpenCL "read_image" Returns a Blank Pixel

Question

Description of Issue and Examples

I am working with OpenCL and I am trying to get an image resize working. I've gone through a few tutorials and I've pretty much ended on trying to write my own kernel for this.

My issue is I cannot get OpenCL to read the pixel from an image that I transfer to it.

Here is Kernel in OpenCL:

const sampler_t SMPL_PREF = CLK_NORMALIZED_COORDS_TRUE |
                            CLK_ADDRESS_CLAMP |
                            CLK_FILTER_LINEAR;

__kernel void image_scaling( __read_only  image2d_t src_img,
                             __write_only image2d_t dest_img,
                                              float WRATIO,
                                              float HRATIO )
{
    int2       coor = (int2)( get_global_id(0), get_global_id(1) );

    float2 origCoor = (float2)(convert_float2(coor).x * WRATIO,
                               convert_float2(coor).y * HRATIO );

    uint4     color = read_imageui( src_img, SMPL_PREF, origCoor );

    // Test write_imageui -- THIS WORKS!!! :(
    //color = (uint4)(20, 255, 20, 255);

    write_imageui( dest_img, coor, color );
}

I am using OpenCL.Net, but I don't think that is what is causing the issue. This is how I am transferring the image from the host to device memory:

ImageFormat clImageFormat = new ImageFormat( ChannelOrder.RGBA,
                                             ChannelType.Unsigned_Int8 );

// Allocate Memory on Device
IMem inputImage2DBuffer = Cl.CreateImage2D( env.Context,
                                            MemFlags.AllocHostPtr |
                                                MemFlags.ReadOnly,
                                            clImageFormat,
                                   (IntPtr) originalWidth,
                                   (IntPtr) originalHeight,
                                   (IntPtr) 0,
                                   (IntPtr) 0, //imageData,
                                        out error );

// Copy the host data to the device
error = Cl.EnqueueWriteImage( env.CommandQueues[ 0 ],  // OpenCL Command Queue
                              inputImage2DBuffer,      // Ptr on device for image
                              Bool.True,               // Blocking write
                              originPtr,               // origin x,y (0,0)
                              regionPtrInput,          // dimension w,h
                     (IntPtr) 0,                       // rowPitch
                     (IntPtr) 0,                       // rowSlice
                              imageData,               // Ptr to byte[] host data
                              0,                       // Number of wait events
                              null,                    // array of wait events
                          out clEvent );

Cl.WaitForEvents( (uint) 1, new Event[] { clEvent } );
CheckError( error, "EnqueueWriteImage" );
clEvent.Dispose( );

I Need Help

I would love if someone could suggest a method of debugging if I am not transferring images properly to the device or if it is something with my device. I am kind of at a dead end. Like I said I dont think it is an issue with OpenCL.Net, but I am almost willing to try anything now. I think that I am probably doing something stupid when transferring the host memory, but I've looked that code over 100 times and rewritten it almost a dozen. Nothing that I can see should be causing this issue.

I know for a fact that I can transfer a regular buffer to the device and read it back, reading images are my only issue. If you look at my OpenCL code, you will see a commented line. By uncommenting it, I can get data back to my host (a lime green color). This leads be to believe the issue is with read_image.

I am working on a 2013 MacBook Air using the Intel 5000 GPU to run my OpenCL code.

Answer 1

Depending on what values you are passing for WRATIO and HRATIO, the problem is very like that you're using CLK_NORMALIZED_COORDS_TRUE and should likely be using CLK_NORMALIZED_COORDS_FALSE. The latter lets you use pixel coordinates instead of 0.0 to 1.0 "normalized" coordinates.

Also, while you're at it, be aware that with CLK_FILTER_LINEAR and float2 coordinates, integral coordinates will give filtered results. Add (float2)(0.5f, 0.5f) to get unfiltered values. In other words, for the pixel at (2,6) read (2.5, 6.5). This will matter when your ratio is near 1.0 so you don't get a fuzzy image.

Answer 2

So, with the help of Dithermaster (thank you!), I've figure out the issue. The issue was that I was not using normalized coordinates when reading back pixels. Here is my current kernel, which is working amazingly!

---

Resize.cl

const sampler_t SMPL_PREF = CLK_NORMALIZED_COORDS_TRUE |
                            CLK_FILTER_LINEAR |
                            CLK_ADDRESS_NONE ;

__kernel void image_scaling( __read_only  image2d_t src_img,
                             __write_only image2d_t dest_img,
                                               uint WIDTH,     // resized width
                                               uint HEIGHT )   // resized height
{
    int2       coor = (int2)( get_global_id(0), get_global_id(1) );
    float2 normCoor = convert_float2(coor) / (float2)( WIDTH, HEIGHT );

    float4    color = read_imagef( src_img, SMPL_PREF, normCoor );

    write_imagef( dest_img, coor, color );
}

---

Again, thanks to Dithermaster for pointing out the spec sheet. OpenCL is very specific when it comes to the filter flags and which read_image you are using. I am super glad I have this fixed finally :)