//
//  subpixel_edge.cpp
//  subpixel_edge
//
//  Created by Mehdi Stapleton on 2015-05-18.
//
//

#include "subpixel_edge.h"

//-------------------------------------------------------------------------
//                 MAIN FUNCTION DEFINITIONS
//-------------------------------------------------------------------------

std::string GetPlatformName (cl_platform_id id)
{
	size_t size = 0;
	clGetPlatformInfo (id, CL_PLATFORM_NAME, 0, NULL, &size);
    
	std::string result;
	result.resize (size);
	clGetPlatformInfo (id, CL_PLATFORM_NAME, size,
                       const_cast<char*> (result.data ()), NULL);
    
	return result;
}

std::string GetDeviceName (cl_device_id id)
{
	size_t size = 0;
	clGetDeviceInfo (id, CL_DEVICE_NAME, 0, NULL, &size);
    
	std::string result;
	result.resize (size);
	clGetDeviceInfo (id, CL_DEVICE_NAME, size,
                     const_cast<char*> (result.data ()), NULL);
    
	return result;
}

void CheckError (cl_int error)
{
	if (error != CL_SUCCESS) {
		std::cerr << "OpenCL call failed with error " << error << std::endl;
		std::exit (1);
	}
}

std::string LoadKernel (const char* name)
{
	std::ifstream in (name);
	std::string result (
                        (std::istreambuf_iterator<char> (in)),
                        std::istreambuf_iterator<char> ());
	return result;
}

cl_program CreateProgram (const std::string& source,
                          cl_context context)
{
	size_t lengths [1] = { source.size () };
	const char* sources [1] = { source.data () };
    
	cl_int error = 0;
	cl_program program = clCreateProgramWithSource (context, 1, sources, lengths, &error);
	CheckError (error);
    
	return program;
}

//-------------------------------------------------------------------------
//         SUBPIXEL EDGE DETECTION (GPU-ACCELERATED)
//-------------------------------------------------------------------------

#ifndef XCODE_COMPILATION

/*\brief apply subpixel edge detection (Devernay1995) on available GPU devices
  1) xGradImg: x-component (column-wise) of gradient
  2) yGradImg: y-component (row-wise) of gradient 
 */
void subpixel_edge_gpu(Image<grad_t> &xGradImgOcl, Image<grad_t> &yGradImgOcl,
                       Image<grad_t> &edgeImgOcl, Image<subpixel_t> &xDeltaImgOcl, Image<subpixel_t> &yDeltaImgOcl)
{
    cl_int error = CL_SUCCESS;
    
    cl_uint platformIdCount = 0;
	clGetPlatformIDs(0, NULL, &platformIdCount);
    
	if (platformIdCount == 0) {
		LOG_ERROR("No OpenCL platform found");
	} else {
		LOG_DEBUG("Found " << platformIdCount << " platform(s)");
	}
    
	std::vector<cl_platform_id> platformIds (platformIdCount);
	clGetPlatformIDs(platformIdCount, platformIds.data (), NULL);
    
	for (cl_uint i = 0; i < platformIdCount; ++i) {
		LOG_DEBUG("\t (" << (i+1) << ") : " << GetPlatformName (platformIds [i]));
	}
    
	cl_uint deviceIdCount = 0;
	clGetDeviceIDs (platformIds [0], CL_DEVICE_TYPE_ALL, 0, NULL,
                    &deviceIdCount);
    
	if (deviceIdCount == 0) {
		LOG_ERROR("No OpenCL devices found");
	} else {
		LOG_DEBUG("Found " << deviceIdCount << " device(s)");
	}
    
	std::vector<cl_device_id> deviceIds (deviceIdCount);
	clGetDeviceIDs (platformIds [0], CL_DEVICE_TYPE_ALL, deviceIdCount,
                    deviceIds.data (), NULL);

#ifdef SHOW_DEBUG_INFO
    cl_uint nMaxWorkItemDims;
    size_t * maxWorkItemSizes;
    size_t nRetWorkItemSizes;
    cl_ulong globalMemSize, localMemSize;
    
	for (cl_uint i = 0; i < deviceIdCount; ++i)
    {
		LOG_DEBUG("\t (" << (i+1) << ") : " << GetDeviceName (deviceIds [i]) );
        
        error = clGetDeviceInfo (deviceIds[i], CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,
                                 sizeof(cl_uint),
                                 &nMaxWorkItemDims,
                                 NULL);
    
        CheckError(error);
        
        maxWorkItemSizes = (size_t*) malloc(sizeof(size_t)*nMaxWorkItemDims);
        
        error = clGetDeviceInfo (deviceIds[i], CL_DEVICE_MAX_WORK_ITEM_SIZES,
                                 sizeof(size_t)*nMaxWorkItemDims,
                                 maxWorkItemSizes,
                                 &nRetWorkItemSizes);
        
        CheckError(error);
        
        error = clGetDeviceInfo (deviceIds[i], CL_DEVICE_GLOBAL_MEM_SIZE,
                                 sizeof(cl_ulong),
                                 &globalMemSize,
                                 NULL);
        
        CheckError(error);
        
        error = clGetDeviceInfo (deviceIds[i], CL_DEVICE_LOCAL_MEM_SIZE,
                                 sizeof(cl_ulong),
                                 &localMemSize,
                                 NULL);
        
        CheckError(error);
        
        LOG_DEBUG("\t Maximum work-item sizes for each dimension: ");
        for( cl_uint j = 0; j<nRetWorkItemSizes/sizeof(size_t); ++j)
        {
            LOG_DEBUG("\t [" << j << "] " << maxWorkItemSizes[j] );
        }
        
        LOG_DEBUG("\t Memory Available (bytes): ");
        LOG_DEBUG("\t Global:" << globalMemSize );
        LOG_DEBUG("\t Local:" << localMemSize );
        
    }
#endif
    
	const cl_context_properties contextProperties [] =
	{
		CL_CONTEXT_PLATFORM, reinterpret_cast<cl_context_properties> (platformIds [0]),
		0, 0
	};
    
	
	cl_context context = clCreateContext (contextProperties, deviceIdCount,
                                          deviceIds.data (), NULL, NULL, &error);
	CheckError (error);
    
	LOG_DEBUG("Context created");
    
	// Create a program from source
    std::string kernelPath = PROJECT_RELPATH("subpixel_edge.cl");
	cl_program program = CreateProgram (LoadKernel (kernelPath.data()),
                                        context);
    
	CheckError ( clBuildProgram (program, deviceIdCount, deviceIds.data (),
                 BUILD_PROGRAM_DEFINE, NULL, NULL));
    
#ifdef SHOW_DEBUG_INFO
    char programBuildLog[BUILD_LOG_BUFFER_SIZE];
    for (cl_uint i = 0; i < deviceIdCount; ++i) {
        CheckError (clGetProgramBuildInfo(program,deviceIds[i],CL_PROGRAM_BUILD_LOG,
                                          BUILD_LOG_BUFFER_SIZE,programBuildLog,NULL));
        LOG_DEBUG(GetDeviceName (deviceIds [i]) << " Build Log: " << programBuildLog);
    }
#endif
    
    // Create kernel
    cl_kernel kernel = clCreateKernel (program, "subpixel_edge", &error);
	CheckError (error);
    
#ifdef SHOW_DEBUG_INFO
    // Display supported image formats
    cl_mem_flags flags;
    cl_uint numMaxSupportedFormats = 50;
    cl_uint numSupportedFormats;
    cl_image_format supportedImgFormats[50];
    clGetSupportedImageFormats(context,
                               flags,
                               CL_MEM_OBJECT_IMAGE2D,
                               numMaxSupportedFormats,
                               supportedImgFormats,
                               &numSupportedFormats);
    
    LOG_DEBUG("Supported Image Formats:");
    for(cl_uint i = 0; i < numSupportedFormats; i++)
    {
        LOG_DEBUG(i << ")" << "channel order: "  << std::hex << std::showbase << supportedImgFormats[i].image_channel_order
                           << ", channel type: " << supportedImgFormats[i].image_channel_data_type);
    }
#endif
    
    //common input image format/descriptions
    std::size_t nColsIn = xGradImgOcl.width;
    std::size_t nRowsIn = xGradImgOcl.height;
	
    cl_mem xGradImage = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR
                                       ,sizeof(grad_t)*nRowsIn*nColsIn
                                       ,xGradImgOcl.pixels.data()
                                       ,&error);
    
    CheckError (error);
    
    cl_mem yGradImage = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR
                                      ,sizeof(grad_t)*nRowsIn*nColsIn
                                      ,yGradImgOcl.pixels.data()
                                      ,&error);
    
    CheckError (error);
    
    //common output image format/descriptions
    int nColsOut = (int)edgeImgOcl.width;
    int nRowsOut = (int)edgeImgOcl.height;
    
    cl_mem edgeImage = clCreateBuffer(context, CL_MEM_WRITE_ONLY
                                     ,sizeof(grad_t)*nRowsOut*nColsOut
                                     ,NULL
                                     ,&error);
    
    CheckError (error);
    
    cl_mem xDeltaImage = clCreateBuffer(context, CL_MEM_WRITE_ONLY
                                        ,sizeof(subpixel_t)*nRowsOut*nColsOut
                                        ,NULL
                                        ,&error);

    CheckError (error);
    
    cl_mem yDeltaImage = clCreateBuffer(context, CL_MEM_WRITE_ONLY
                                        ,sizeof(subpixel_t)*nRowsOut*nColsOut
                                        ,NULL
                                        ,&error);
    
    CheckError (error);
    
	cl_command_queue queue = clCreateCommandQueue (context, deviceIds [1],
                                                   0, &error);
	CheckError (error);
    
    // Setup the kernel arguments
	clSetKernelArg (kernel, 0, sizeof (cl_mem), &xGradImage);
    clSetKernelArg (kernel, 1, sizeof (cl_mem), &yGradImage);
    clSetKernelArg (kernel, 2, sizeof (cl_mem), &edgeImage);
	clSetKernelArg (kernel, 3, sizeof (cl_mem), &xDeltaImage);
    clSetKernelArg (kernel, 4, sizeof (cl_mem), &yDeltaImage);
    clSetKernelArg (kernel, 5, sizeof (int), &nRowsOut);
    clSetKernelArg (kernel, 6, sizeof (int), &nColsOut);
    
	// Run the processing
	std::size_t offset [2]    = { 0, 0 };
	std::size_t size [2]      = { nColsIn, nRowsIn};
    std::size_t localSize [2] = { LOCAL_WORK_SIZE, LOCAL_WORK_SIZE};
    
    if( nColsIn%LOCAL_WORK_SIZE || nRowsIn%LOCAL_WORK_SIZE )
    {
        CheckError (clEnqueueNDRangeKernel (queue, kernel, 2, offset, size, NULL,
                                            0, NULL, NULL));
    }
    else
    {
        CheckError (clEnqueueNDRangeKernel (queue, kernel, 2, offset, size, localSize,
                                            0, NULL, NULL));
    }
	
	
    CheckError (clFinish(queue));
    
    //Retrieve output images from GPU
    std::size_t bufOffset       = 0;
    std::size_t bufSizeGradImg  = sizeof(grad_t)*nColsOut*nRowsOut;
    std::size_t bufSizeDeltaImg = sizeof(subpixel_t)*nColsOut*nRowsOut;
    
    CheckError (clEnqueueReadBuffer (queue, edgeImage, CL_TRUE,
                                     bufOffset, bufSizeGradImg,
                                     edgeImgOcl.pixels.data (), 0, NULL, NULL));
    
    CheckError (clEnqueueReadBuffer (queue, xDeltaImage, CL_TRUE,
                                     bufOffset, bufSizeDeltaImg,
                                     xDeltaImgOcl.pixels.data (), 0, NULL, NULL));
    
    CheckError (clEnqueueReadBuffer (queue, yDeltaImage, CL_TRUE,
                                     bufOffset, bufSizeDeltaImg,
                                     yDeltaImgOcl.pixels.data (), 0, NULL, NULL));
    
    //clean-up
	clReleaseMemObject (xGradImage);
    clReleaseMemObject (yGradImage);
	clReleaseMemObject (edgeImage);
    clReleaseMemObject (xDeltaImage);
    clReleaseMemObject (yDeltaImage);
    
	clReleaseCommandQueue (queue);
	
	clReleaseKernel (kernel);
    
	clReleaseProgram (program);
    
	clReleaseContext (context);
    
    
}//end of subpixel edge gpu

#endif//XCODE_COMPILATION



