How to display an image inside kernel using opencl?

Question

I am new to opencl. The task is:

1. Load an pre-existing image
2. Write Host code using opencl to send the image ptr to kernel
3. Calculate hsl threshold of the loaded image inside kernel
4. Display the threshold or binary image

I ve used opencv to load a pre-existing 2D image in my program. And I used open cl buffer objects to allocate memory and have send image pointer to the kernel. After kernel execution in order to display the calculated image from the kernel I need clEnqueueReadBuffer. Then I use opencv to display the image from the host. I ve attached code below

As this takes more time on GPU and CPU I thought to switch over to image memory.

But I like to know whether usage of images also need clenqueueReadImage to copy image from kernel to host or do we any way to display the threshold image in kernel itself?

//My code using opencl buffers    
IplImage *src = cvLoadImage("../Input/im2.png",CV_LOAD_IMAGE_COLOR );

int a=src->height;
int b=src->width;

cl_context CreateContext()
{
    cl_int errNum;
    cl_uint numPlatforms;
    cl_platform_id firstPlatformId;
    cl_context context = NULL;
    errNum = clGetPlatformIDs(1, &firstPlatformId, &numPlatforms);
    if (errNum != CL_SUCCESS || numPlatforms <= 0)
    {
        std::cerr << "Failed to find any OpenCL platforms." << std::endl;
        return NULL;
    }

    cl_context_properties contextProperties[] =
    {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)firstPlatformId,
        0
    };
    context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU,
                                      NULL, NULL, &errNum);
    if (errNum != CL_SUCCESS)
    {
        std::cout << "Could not create GPU context, trying CPU..." << std::endl;
        context = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_CPU, NULL, NULL, &errNum);
        if (errNum != CL_SUCCESS)
        {
            std::cerr << "Failed to create an OpenCL GPU or CPU context." << std::endl;
            return NULL;
        }
    }
    return context;
}


cl_command_queue CreateCommandQueue(cl_context context, cl_device_id *device)
{
    cl_int errNum;
    cl_device_id *devices;
    cl_command_queue commandQueue = NULL;
    size_t deviceBufferSize = -1;
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &deviceBufferSize);
    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Failed call to clGetContextInfo(...,GL_CONTEXT_DEVICES,...)";
        return NULL;
    }

    if (deviceBufferSize <= 0)
    {
        std::cerr << "No devices available.";
        return NULL;
    }
    devices = new cl_device_id[deviceBufferSize / sizeof(cl_device_id)];
    errNum = clGetContextInfo(context, CL_CONTEXT_DEVICES, deviceBufferSize, devices, NULL);
    if (errNum != CL_SUCCESS)
    {
        delete [] devices;
        std::cerr << "Failed to get device IDs";
        return NULL;
    }

    commandQueue = clCreateCommandQueue(context, devices[0],CL_QUEUE_PROFILING_ENABLE, &errNum );

    if (commandQueue == NULL)
    {
        delete [] devices;
        std::cerr << "Failed to create commandQueue for device 0";
        return NULL;
    }

    *device = devices[0];
    delete [] devices;
    return commandQueue;
}

cl_program CreateProgram(cl_context context, cl_device_id device, const char* fileName)
{
    cl_int errNum;
    cl_program program;

    std::ifstream kernelFile(fileName, std::ios::in);
    if (!kernelFile.is_open())
    {
        std::cerr << "Failed to open file for reading: " << fileName << std::endl;
        return NULL;
    }

    std::ostringstream oss;
    oss << kernelFile.rdbuf();

    std::string srcStdStr = oss.str();
    const char *srcStr = srcStdStr.c_str();
    program = clCreateProgramWithSource(context, 1,
                                        (const char**)&srcStr,
                                        NULL, NULL);
    if (program == NULL)
    {
        std::cerr << "Failed to create CL program from source." << std::endl;
        return NULL;
    }
    errNum = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);

    if (errNum != CL_SUCCESS)
    {
        char buildLog[16384];
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG,
                              sizeof(buildLog), buildLog, NULL);
        std::cerr << "Error in kernel: " << std::endl;
        std::cerr << buildLog;
        clReleaseProgram(program);
        return NULL;
    }
    return program;
}

bool CreateMemObjects(cl_context context, cl_mem memObjects[2], unsigned char *src_ptr)
{
    memObjects[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(unsigned char) *(a*b*3) , src_ptr , NULL);
    memObjects[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(unsigned char) *(a*b) , NULL, NULL);

   if (memObjects[0] == NULL || memObjects[1] == NULL)
    {
        std::cerr << "Error creating memory objects" << std::endl;
        return false;
    }
    return true;
}

void Cleanup(cl_context context, cl_command_queue commandQueue, cl_program program, cl_kernel kernel, cl_mem memObjects[2])
{
    for (int i = 0; i < 2; i++)
    {
        if (memObjects[i] != 0)
            clReleaseMemObject(memObjects[i]);
    }
    if (commandQueue != 0)
        clReleaseCommandQueue(commandQueue);

    if (kernel != 0)
        clReleaseKernel(kernel);

    if (program != 0)
        clReleaseProgram(program);

    if (context != 0)
        clReleaseContext(context);
}


int main()
{
    cl_context context = 0;
    cl_command_queue commandQueue = 0;
    cl_program program = 0;
    cl_device_id device = 0;
    cl_kernel kernel = 0;
    cl_mem memObjects[2] = { 0,0 };
    cl_int errNum;
    cl_event myEvent;

    cl_ulong start_time,end_time;
    double kernelExecTimeNs;

    IplImage *thres_img1 = cvCreateImage(cvGetSize(src), IPL_DEPTH_8U, 1);

    unsigned char *tur_image1,*src_ptr;
    tur_image1 = (unsigned char*) malloc((a*b) * sizeof(unsigned char));
    src_ptr = (unsigned char*) malloc ((a*b*3) * sizeof(unsigned char));

    context = CreateContext();
    if (context == NULL)
    {
        std::cerr << "Failed to create OpenCL context." <<std::endl;
        return 1;
    }

    commandQueue = CreateCommandQueue(context, &device);
    if (commandQueue == NULL)
    {
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

    program = CreateProgram(context, device, "hsl_threshold.cl");
    if (program == NULL)
    {
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

    kernel = clCreateKernel(program, "HSL_threshold", NULL);
    if (kernel == NULL)
    {
        std::cerr << "Failed to create kernel" << std::endl;
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

    printf("height:%d\n",a);//image height
    printf("width:%d\n",b);//image width

    cvShowImage("color image",src);
    cvWaitKey(0);

    memcpy(src_ptr,src->imageData,(a*b*3));

    if (!CreateMemObjects(context, memObjects, src_ptr))
    {
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

        errNum  = clSetKernelArg(kernel, 0, sizeof(cl_mem), &memObjects[0]);
    errNum |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &memObjects[1]);

    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Error setting kernel arguments" << std::endl;
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

    cout<<"Kernel arguments set successfully";
    size_t globalWorkSize[1]={a*b};
    size_t localWorkSize[1]={512};

    errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL, globalWorkSize, localWorkSize, 0, NULL, &myEvent);

    clWaitForEvents(1,&myEvent);

    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Error queuing kernel for execution." << std::endl;
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }
    clFinish(commandQueue);


    clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_START, sizeof(start_time), &start_time, NULL);
    clGetEventProfilingInfo(myEvent, CL_PROFILING_COMMAND_END, sizeof(end_time), &end_time, NULL);

    kernelExecTimeNs = end_time-start_time;

    printf("\nExecution time in milliseconds = %0.3f ms\n",( kernelExecTimeNs / 1000000.0) );
    cout<<"\n Kernel timings \n"<<kernelExecTimeNs<<"seconds";

    errNum = clEnqueueReadBuffer(commandQueue, memObjects[1], CL_TRUE,
                                 0, (a*b) * sizeof(unsigned char), tur_image1,
                                 0, NULL, NULL);

    if (errNum != CL_SUCCESS)
    {
        std::cerr << "Error reading result buffer." << std::endl;
        Cleanup(context, commandQueue, program, kernel, memObjects);
        return 1;
    }

    memcpy(thres_img1->imageData,tur_image1,sizeof(unsigned char)*(a*b));

    cvShowImage( "hsl_thresh",thres_img1);
    cvSaveImage( "../Output/hsl_threshold.png",thres_img1);
    cvWaitKey(0);

    std::cout<<std::endl;
    std::cout<<"Image displayed Successfully"<<std::endl;

    Cleanup(context,commandQueue,program,kernel,memObjects);
    printf("\n Free opencl resources");
    std::cin.get();
    return 0;

}

Answer 1

There are ways to directly process data calculated by OpenCL via OpenGL. Your OCL implementation must support the extension cl_khr_gl_sharing.
This mode is called CL/GL-Interop Mode.

If you create an OpenGL-instance first and initialise OpenCL with the pointers to your GL-instance, it is possible for each implementation to access each others data.

(All snippets are taken from code using CL-C++-Bindings, I guess it is okay for the general understanding)

cl_context_properties properties[] = 
  // Take this line to create an OCL context in GL-CL-interop-mode. 
  // OpenGL must already be initialised. 
  // For interop init see: http://www.khronos.org/registry/cl/extensions/khr/cl_khr_gl_sharing.txt
  // USING: CL_GL_CONTEXT_KHR: Rendering Context [Use your OGL-HGLRC variable or do wglGetCurrentContext(); ]
  //   AND: CL_WGL_HDC_KHR: Device Context [Use your OGL-HDC variable or do wglGetCurrentDC(); ]
  { 
    CL_CONTEXT_PLATFORM, (cl_context_properties)(_platforms->at(0))(), 
    CL_GL_CONTEXT_KHR, (cl_context_properties)myGL->hRC, 
    CL_WGL_HDC_KHR, (cl_context_properties)myGL->hDC, 0
  };

Now you can create OCL-images based on OGL textures

//The following data can be accessed both from OCL and OGL
cl::Image2D imageFromGL = new cl::Image2DGL(*_context, CL_MEM_READ_WRITE,  GL_TEXTURE_2D, 0, myGL->textures[0]);

Before using the memory in OCL, you have to ask OGL to release it

//Ask OGL to release memory. All OGL actions must be finished before doing so!
_queue->enqueueAcquireGLObjects(&imageFromGL, NULL, &evt);

Now, do what you want, then give it back to OGL:

//Hand memory back to OGL. All OCL actions must be finished before doing so!
_queue->enqueueReleaseGLObjects(&imageFromGL, NULL, &evt);

And finally you can use OpenGL code to display the data on the screen.