#include <stdlib.h>
#include <stdio.h>
#include <string.h>
#include <CL/opencl.h>

// Utility function to get file size
int getFileSize(FILE *fileArg)
{
    if (fseek(fileArg, 0, SEEK_END))
        return -1;
    int fileSize = ftell(fileArg);
    if (fseek(fileArg, 0, SEEK_SET))
        return -1;
    return fileSize;
}

// Utility function to get Intel OpenCL platform
cl_platform_id choosePlatform()
{
    // Enumerate compute platforms on board
    cl_platform_id platformList[4];
    cl_uint numPlatforms = 0;
    cl_int status = clGetPlatformIDs(4, platformList, &numPlatforms);
    if (status != CL_SUCCESS || numPlatforms < 1)
    {
        printf("CL Platform list query error\n");
        return NULL;
    }

    // Choose Intel's one
    cl_platform_id clPlatform = NULL;
    for (cl_uint i = 0; i < numPlatforms && clPlatform == NULL; i++)
    {
        size_t nameStrSize = 0;
        char *nameStr = NULL;
        status = clGetPlatformInfo(platformList[i], CL_PLATFORM_NAME, 0, NULL, &nameStrSize);
        if (status == CL_SUCCESS)
        {
            nameStr = (char*)malloc(sizeof(char) * nameStrSize);
            status = clGetPlatformInfo(platformList[i], CL_PLATFORM_NAME, sizeof(char) * nameStrSize, nameStr, NULL);
        }
        if (status == CL_SUCCESS && strstr(nameStr, "Intel") != NULL)
            clPlatform = platformList[i];
        if (nameStr != NULL)
            free(nameStr);
    }

    return clPlatform;
}

cl_program buildProgram(cl_context clContext, cl_device_id clDevice)
{
    // Load OpenCL source from file
    FILE *srcFile = fopen("kernels.cl", "rb");
    if (srcFile == NULL)
    {
        printf("Source file open error\n");
        return NULL;
    }
    size_t srcSize = getFileSize(srcFile);
    if (srcSize <= 0)
    {
        printf("File size error\n");
        return NULL;
    }

    // Read program source to string
    char *programSrc = (char*)malloc((srcSize + 1) * sizeof(char));
    size_t readSize = fread(programSrc, sizeof(char), srcSize, srcFile);
    fclose(srcFile);
    if (readSize < srcSize)
    {
        free(programSrc);
        printf("Read program source file error\n");
        return NULL;
    }
    // Append null-character (it may be missed in source file)
    programSrc[srcSize] = '\n';
    srcSize++;

    // Create program with source
    cl_program clProgram = clCreateProgramWithSource(clContext, 1, (const char**)&programSrc, &srcSize, NULL);
    free(programSrc);
    cl_int buildStatus = !CL_SUCCESS;
    if (clProgram != NULL)
        buildStatus = clBuildProgram(clProgram, 1, &clDevice, "", NULL, NULL);
    if (buildStatus != CL_SUCCESS && clProgram != NULL)
    {
        clReleaseProgram(clProgram);
        clProgram = NULL;
    }

    return clProgram;
}

typedef struct _ComputeHandle
{
    cl_kernel kernel;
    cl_mem outputBuffer;
    cl_mem internalBuffer;
    size_t globalNDRange[2];
    size_t localNDRange[2];

} ComputeHandle;


cl_int initializeComputeHandle(ComputeHandle *handle,
                               cl_context context,
                               cl_program program,
                               cl_mem *inputBuffer,
                               unsigned int bufferWidth,
                               unsigned int bufferHeight)
{
    handle->kernel = clCreateKernel(program, "kernel0", NULL);
    if (handle->kernel == NULL)
        return !CL_SUCCESS;

    handle->outputBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                          bufferWidth * bufferHeight * sizeof(cl_uchar),
                                          NULL, NULL);
    if (handle->outputBuffer == NULL)
        return !CL_SUCCESS;

    cl_uint iterationsInnerCount = 2;
    size_t internalBufferSize = (2 * iterationsInnerCount + 1) * sizeof(cl_float4);
    handle->internalBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY,
                                            internalBufferSize,
                                            NULL, NULL);
    if (handle->internalBuffer == NULL)
        return !CL_SUCCESS;

    cl_uint iterationsCount = 5;
    handle->globalNDRange[0] = bufferWidth / (4 * iterationsCount);
    handle->globalNDRange[1] = bufferHeight;
    handle->localNDRange[0] = 16;
    handle->localNDRange[1] = 32;

    cl_uint kernelArgNum = 0;

    clSetKernelArg(handle->kernel, kernelArgNum++,
                   sizeof(cl_mem), inputBuffer);
    clSetKernelArg(handle->kernel, kernelArgNum++,
                   sizeof(cl_mem), &handle->outputBuffer);

    cl_uint bufferWidthChar4 = bufferWidth / 4;
    clSetKernelArg(handle->kernel, kernelArgNum++,
                   sizeof(cl_uint), &bufferWidthChar4);

    clSetKernelArg(handle->kernel, kernelArgNum++,
                   sizeof(cl_mem), &handle->internalBuffer);

    clSetKernelArg(handle->kernel, kernelArgNum++,
                   sizeof(cl_uint), &iterationsCount);
    clSetKernelArg(handle->kernel, kernelArgNum++,
                   sizeof(cl_uint), &iterationsInnerCount);

    size_t localMemSize = handle->localNDRange[0] * handle->localNDRange[1] *
                          sizeof(cl_uchar4) * 2;
    clSetKernelArg(handle->kernel, kernelArgNum++, localMemSize, NULL);

    return CL_SUCCESS;
}

void releaseComputeHandle(ComputeHandle *handle)
{
    if (handle->outputBuffer != NULL)
        clReleaseMemObject(handle->outputBuffer);
    handle->outputBuffer = NULL;

    if (handle->internalBuffer != NULL)
        clReleaseMemObject(handle->internalBuffer);
    handle->internalBuffer = NULL;

    if (handle->kernel != NULL)
        clReleaseKernel(handle->kernel);
    handle->kernel = NULL;
}


cl_int enqueueComputeWork(const ComputeHandle *handle,
                          cl_command_queue queue,
                          cl_uint numWaitEvents,
                          const cl_event *eventWaitList,
                          cl_event *event)
{
    cl_int status = clEnqueueNDRangeKernel(
                queue, handle->kernel,
                2, 0, handle->globalNDRange, handle->localNDRange,
                numWaitEvents, eventWaitList, event);
    return status;
}


int main()
{
    cl_int status = CL_SUCCESS;

    // Get OpenCL platform
    cl_platform_id clPlatform = choosePlatform();
    if (clPlatform == NULL)
    {
        printf("No Intel Platform onboard\n");
        return -1;
    }

    // Create OpenCL context
    cl_context_properties contextProperties[] = {CL_CONTEXT_PLATFORM, (cl_context_properties)clPlatform, 0};
    cl_context clContext = clCreateContextFromType(contextProperties, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL);
    if (clContext == NULL)
    {
        printf("Context create error\n");
        return -2;
    }

    // Get GPU device
    cl_device_id clDevice = NULL;
    status = clGetContextInfo(clContext, CL_CONTEXT_DEVICES, sizeof(cl_device_id), &clDevice, NULL);
    if (status != CL_SUCCESS)
    {
        clReleaseContext(clContext);
        printf("Get context info error\n");
        return -3;
    }

    // Build CL Program
    cl_program clProgram = buildProgram(clContext, clDevice);
    if (clProgram == NULL)
    {
        clReleaseContext(clContext);
        printf("Program build error\n");
        return -4;
    }

    // Create command queue
    cl_command_queue clQueue = NULL;
    if (clProgram != NULL)
        clQueue = clCreateCommandQueue(clContext, clDevice, 0, NULL);


    // Initialize input buffer
    const unsigned int bufferWidth = 1280;
    const unsigned int bufferHeight = 704;
    cl_mem inputBuffer = NULL;
    if (clQueue != NULL)
        inputBuffer = clCreateBuffer(clContext, CL_MEM_READ_ONLY,
                                     bufferWidth * bufferWidth * sizeof(cl_uchar),
                                     NULL, NULL);
    if (inputBuffer == NULL)
    {
        if (clQueue != NULL)
            clReleaseCommandQueue(clQueue);
        clReleaseProgram(clProgram);
        clReleaseContext(clContext);
        printf("OpenCL resources allocation failed\n");
        return -8;
    }

    // Create handles
    const unsigned int handleCount = 7;
    ComputeHandle firstStepHandles[handleCount];
    ComputeHandle secondStepHandles0[handleCount];
    ComputeHandle secondStepHandles1[handleCount];
    ComputeHandle secondStepHandles2[handleCount];
    ComputeHandle thirdStepHandles0[handleCount];
    ComputeHandle thirdStepHandles1[handleCount];
    ComputeHandle thirdStepHandles2[handleCount];

    memset(firstStepHandles, 0, handleCount * sizeof(ComputeHandle));
    memset(secondStepHandles0, 0, handleCount * sizeof(ComputeHandle));
    memset(secondStepHandles1, 0, handleCount * sizeof(ComputeHandle));
    memset(secondStepHandles2, 0, handleCount * sizeof(ComputeHandle));
    memset(thirdStepHandles0, 0, handleCount * sizeof(ComputeHandle));
    memset(thirdStepHandles1, 0, handleCount * sizeof(ComputeHandle));
    memset(thirdStepHandles2, 0, handleCount * sizeof(ComputeHandle));


    for (unsigned int handleIndex = 0; handleIndex < handleCount &&
                                       status == CL_SUCCESS; handleIndex++)
    {
        status = initializeComputeHandle(&firstStepHandles[handleIndex], clContext, clProgram,
                                         &inputBuffer, bufferWidth, bufferHeight);
        if (status != CL_SUCCESS)
            continue;

        status = initializeComputeHandle(&secondStepHandles0[handleIndex], clContext, clProgram,
                                         &firstStepHandles[handleIndex].outputBuffer, bufferWidth, bufferHeight);
        if (status != CL_SUCCESS)
            continue;

        status = initializeComputeHandle(&secondStepHandles1[handleIndex], clContext, clProgram,
                                         &firstStepHandles[handleIndex].outputBuffer, bufferWidth, bufferHeight);
        if (status != CL_SUCCESS)
            continue;

        status = initializeComputeHandle(&secondStepHandles2[handleIndex], clContext, clProgram,
                                         &firstStepHandles[handleIndex].outputBuffer, bufferWidth, bufferHeight);
        if (status != CL_SUCCESS)
            continue;

        status = initializeComputeHandle(&thirdStepHandles0[handleIndex], clContext, clProgram,
                                         &firstStepHandles[handleIndex].outputBuffer, bufferWidth, bufferHeight);
        if (status != CL_SUCCESS)
            continue;

        status = initializeComputeHandle(&thirdStepHandles1[handleIndex], clContext, clProgram,
                                         &firstStepHandles[handleIndex].outputBuffer, bufferWidth, bufferHeight);
        if (status != CL_SUCCESS)
            continue;

        status = initializeComputeHandle(&thirdStepHandles2[handleIndex], clContext, clProgram,
                                         &firstStepHandles[handleIndex].outputBuffer, bufferWidth, bufferHeight);
        if (status != CL_SUCCESS)
            continue;
    }

    if (status != CL_SUCCESS)
    {
        for (unsigned int handleIndex = 0; handleIndex < handleCount; handleIndex++)
        {
            releaseComputeHandle(&firstStepHandles[handleIndex]);
            releaseComputeHandle(&secondStepHandles0[handleIndex]);
            releaseComputeHandle(&secondStepHandles1[handleIndex]);
            releaseComputeHandle(&secondStepHandles2[handleIndex]);
            releaseComputeHandle(&thirdStepHandles0[handleIndex]);
            releaseComputeHandle(&thirdStepHandles1[handleIndex]);
            releaseComputeHandle(&thirdStepHandles2[handleIndex]);
        }
        clReleaseMemObject(inputBuffer);
        clReleaseCommandQueue(clQueue);
        clReleaseProgram(clProgram);
        clReleaseContext(clContext);
        printf("Compute resources initialization failed\n");
        return -8;
    }

    cl_event firstStepEvents[handleCount];
    cl_event secondStepEvents[handleCount];
    cl_event thirdStepEvents[handleCount];

    for (unsigned int iteration = 0; iteration < 5 && status == CL_SUCCESS; iteration++)
    {
        for (unsigned int handleIndex = 0; handleIndex < handleCount &&
                                           status == CL_SUCCESS; handleIndex++)
        {
            status = enqueueComputeWork(&firstStepHandles[handleIndex], clQueue,
                                        0, NULL, &firstStepEvents[handleIndex]);
        }

        for (unsigned int handleIndex = 0; handleIndex < handleCount &&
                                           status == CL_SUCCESS; handleIndex++)
        {
            cl_event secondEventsTemp[3] = {NULL, NULL, NULL};
            if (status == CL_SUCCESS)
                status = enqueueComputeWork(&secondStepHandles0[handleIndex], clQueue,
                                            1, &firstStepEvents[handleIndex], &secondEventsTemp[0]);
            if (status == CL_SUCCESS)
                status = enqueueComputeWork(&secondStepHandles1[handleIndex], clQueue,
                                            1, &firstStepEvents[handleIndex], &secondEventsTemp[1]);
            if (status == CL_SUCCESS)
                status = enqueueComputeWork(&secondStepHandles2[handleIndex], clQueue,
                                            1, &firstStepEvents[handleIndex], &secondEventsTemp[2]);

            if (firstStepEvents[handleIndex] != NULL)
                clReleaseEvent(firstStepEvents[handleIndex]);

            if (status == CL_SUCCESS)
                clEnqueueMarkerWithWaitList(clQueue, 3, secondEventsTemp, &secondStepEvents[handleIndex]);

            if (secondEventsTemp[0] != NULL)
                clReleaseEvent(secondEventsTemp[0]);
            if (secondEventsTemp[1] != NULL)
                clReleaseEvent(secondEventsTemp[1]);
            if (secondEventsTemp[2] != NULL)
                clReleaseEvent(secondEventsTemp[2]);
        }

        for (unsigned int handleIndex = 0; handleIndex < handleCount &&
                                           status == CL_SUCCESS; handleIndex++)
        {
//            clWaitForEvents(1, &secondStepEvents[handleIndex]);

            cl_event thirdEventsTemp[3] = {NULL, NULL, NULL};
            if (status == CL_SUCCESS)
                status = enqueueComputeWork(&thirdStepHandles0[handleIndex], clQueue,
                                            1, &secondStepEvents[handleIndex], &thirdEventsTemp[0]);
            if (status == CL_SUCCESS)
                status = enqueueComputeWork(&thirdStepHandles1[handleIndex], clQueue,
                                            1, &secondStepEvents[handleIndex], &thirdEventsTemp[1]);
            if (status == CL_SUCCESS)
                status = enqueueComputeWork(&thirdStepHandles2[handleIndex], clQueue,
                                            1, &secondStepEvents[handleIndex], &thirdEventsTemp[2]);

            if (firstStepEvents[handleIndex] != NULL)
                clReleaseEvent(secondStepEvents[handleIndex]);

            if (status == CL_SUCCESS)
                clEnqueueMarkerWithWaitList(clQueue, 3, thirdEventsTemp, &thirdStepEvents[handleIndex]);

            if (thirdEventsTemp[0] != NULL)
                clReleaseEvent(thirdEventsTemp[0]);
            if (thirdEventsTemp[1] != NULL)
                clReleaseEvent(thirdEventsTemp[1]);
            if (thirdEventsTemp[2] != NULL)
                clReleaseEvent(thirdEventsTemp[2]);
        }

        for (unsigned int handleIndex = 0; handleIndex < handleCount &&
                                           status == CL_SUCCESS; handleIndex++)
        {
            if (thirdStepEvents[handleIndex] != NULL)
            {
                clWaitForEvents(1, &thirdStepEvents[handleIndex]);
                clReleaseEvent(thirdStepEvents[handleIndex]);
            }
        }
    }

    if (status != CL_SUCCESS)
        printf("Processing failed\n");
    else
        printf("Processing finished\n");


    // Wait to finish and release resources
    clFinish(clQueue);
    for (unsigned int handleIndex = 0; handleIndex < handleCount; handleIndex++)
    {
        releaseComputeHandle(&firstStepHandles[handleIndex]);
        releaseComputeHandle(&secondStepHandles0[handleIndex]);
        releaseComputeHandle(&secondStepHandles1[handleIndex]);
        releaseComputeHandle(&secondStepHandles2[handleIndex]);
    }
    clReleaseMemObject(inputBuffer);
    clReleaseCommandQueue(clQueue);
    clReleaseProgram(clProgram);
    clReleaseContext(clContext);

    return 0;
}

