OpenCL "read_imageui " always returns zero 0

Question

I have written a simple OpenCL program with an objective to make a copy of input image using OpenCL image2d struct. It seemed like a simple job to do but I have been stuck at it.

The kernel has "read_imageui" which always returns zero value. The input image is a all white jpeg image.

Image loading is done using OpenCV imread.

Here is the Kernel :

const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;


__kernel void copy(__read_only image2d_t in, __write_only image2d_t out)
{
    int idx = get_global_id(0);
    int idy = get_global_id(1);

    int2 pos = (int2)(idx,idy);
    uint4 pix = read_imageui(in,smp,pos);

    write_imageui(out,pos,pix);

}

Here is the host code :

int main(){
    //get all platforms (drivers)
    std::vector<cl::Platform> all_platforms;
    cl::Platform::get(&all_platforms);
    if(all_platforms.size()==0){
        std::cout<<" No platforms found. Check OpenCL installation!\n";
        exit(1);
    }
    cl::Platform default_platform=all_platforms[0];
    std::cout << "Using platform: "<<default_platform.getInfo<CL_PLATFORM_NAME>()<<"\n";
    std::cout <<" Platform Version: "<<default_platform.getInfo<CL_PLATFORM_VERSION>() <<"\n";
    //cout << "Image 2D support : " << default_platform.getInfo<CL_DEVICE_IMAGE_SUPPORT>()<<"\n";

    //get default device of the default platform
    std::vector<cl::Device> all_devices;
    default_platform.getDevices(CL_DEVICE_TYPE_ALL, &all_devices);
    if(all_devices.size()==0){
        std::cout<<" No devices found. Check OpenCL installation!\n";
        exit(1);
    }
    cl::Device default_device=all_devices[0];
    std::cout<< "Using device: "<<default_device.getInfo<CL_DEVICE_NAME>()<<"\n";

    //creating a context

    cl::Context context(default_device);
    //cl::Program::Sources sources;
    //sources.push_back(LoadKernel('kenel2.cl'));


    //load kernel coad
    cl::Program program(context,LoadKernel("image_test.cl"));

    //build kernel code
    if(program.build(all_devices)!=CL_SUCCESS){
        std::cout<<" Error building: "<<program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(default_device)<<"\n";
        exit(1);
    }

   /* IMAGE FORMTS */

    // Determine and show image format support 
    vector<cl::ImageFormat > supportedFormats;
    context.getSupportedImageFormats(CL_MEM_READ_ONLY,CL_MEM_OBJECT_IMAGE2D,&supportedFormats);

    cout <<"No. of supported formats " <<supportedFormats.size()<<endl;
    Mat white = imread("white_small.jpg");

    cvtColor(white, white, CV_BGR2RGBA);

    //white.convertTo(white,CV_8UC4);
    Mat out = Mat(white);
    out.setTo(Scalar(0));

    char * inbuffer = reinterpret_cast<char *>(white.data);
    char * outbuffer = reinterpret_cast<char *>(out.data);

    //cout <<"Type of input : " <<white.type<<endl;

    int sizeOfImage = white.cols * white.rows * white.channels();
    int outImageSize = white.cols * white.rows * white.channels();
    int w = white.cols;
    int h = white.rows;

    cout <<"Creating Images ... "<<endl;
    cout <<"Dimensions ..." <<w << " x "<<h<<endl;


    const cl::ImageFormat format(CL_RGBA, CL_UNSIGNED_INT8);
    cl::Image2D imageSrc(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, format, white.cols, white.rows,0,inbuffer);
    cl::Image2D imageDst(context, CL_MEM_WRITE_ONLY, format , white.cols, white.rows,0,NULL);

    cout <<"Creating Kernel Program ... "<<endl;

    cl::Kernel kernelCopy(program, "copy");
    kernelCopy.setArg(0, imageSrc);
    kernelCopy.setArg(1, imageDst);

    cout <<"Creating Command Queue ... "<<endl;
    cl::CommandQueue queue(context, default_device);

    cout <<"Executing Kernel ... "<<endl;
    int64 e = getTickCount();
    for(int i = 0 ; i < 100 ; i ++)
    {


        queue.enqueueNDRangeKernel(kernelCopy, cl::NullRange, cl::NDRange(w, h), cl::NullRange);
        queue.finish();
    }

    cout <<((getTickCount() - e) / getTickFrequency())/100 <<endl;;

    cl::size_t<3> origin;
    cl::size_t<3> size;
    origin[0] = 0;
    origin[1] = 0;
    origin[2] = 0;
    size[0] = w;
    size[1] = h;
    size[2] = 1;

    cout <<"Transfering Images ... "<<endl;
    //unsigned char *tmp = new unsigned char (w * h * 4);
    //CL_TRUE means that it waits for the entire image to be copied before continuing
    queue.enqueueReadImage(imageDst, CL_TRUE, origin, size, 0, 0, outbuffer);
    queue.finish();
    imwrite("result.jpg",out);
    /* OLD CODE ==================================================*/

    return 0;
}

However if I change the kernel as

uint4 pix2 = (uint4)(255,255,255,1); 
 write_imageui(out,pos,pix2);

It outputs a white image. Which means there is something wrong with how I am using the read_image

Answer 1

When write:

Mat out = Mat(white);

It is like a shallow copy of white to out. Bot white.data and out.data pointers will be pointing to same memory and reference count will be incremented. So, when you call out.setTo, white Mat will also see same change. Declaring out as below might be good idea:

Mat out = Mat(white.size,CV_8UC(white.channels()));

Answer 2

it came out to be something related to "reference counting" on Mat copy constructor.

if instead of using

Mat white = imread("white_small.jpg");

    cvtColor(white, white, CV_BGR2RGBA);

    //white.convertTo(white,CV_8UC4);
    Mat out = Mat(white);

Initialize the output matrix "out" as

Mat out = Mat(white.size,CV_8UC4)

then it works fine.

I couldn't comprehend completely what exactly caused it but I know that it is due to "reference counting" of Mat copy constructor when used as first syntax.