#include <iostream>
#include <memory>
#include <string>
#include "Eigen/Core"

#ifdef HAVE_CL2
    std::cout << "using opencl 2.0" << std::endl;
    #include <CL/cl2.hpp>
#else
    #include <CL/cl.hpp>
#endif

using namespace cl;


int main()
{
    // declarations
    cl_int err = 0;
    const char* raw_src = 
    "__kernel void iMatMult(const int Mdim, const int Ndim, const int Pdim,__global const int *A, __global const int *B, __global int *C) { "
    "	    int k; "
    "	    const int globalRow = get_global_id(0);"
    "	    const int globalCol = get_global_id(1);"
    "	    int tmp = 0;"
    "	    for(k=0; k < Pdim; k++){ "
    "		tmp += A[k*Mdim+globalRow] * B[globalCol*Pdim+k]; "
    "	    } "
    "	    C[globalCol*Mdim+globalRow] = tmp; "
    "	}";

    std::string src = (std::string)(raw_src);
    cl_device_type ocl_device;

    Program::Sources sources;
    sources.push_back({src.c_str(), src.length()});

    Eigen::Matrix4i Am;
    Am << 1, 2, 3, 4,
         5, 6, 7, 8,
         9, 10, 11, 12,
	 13, 14, 15, 16;
    std::cout << Am;

    std::cout << "" << std::endl;

    Eigen::Matrix4i Bm;
    Bm << 1, 2, 3, 4,
         5, 6, 7, 8,
         9, 10, 11, 12,
	 13, 14, 15, 16;
    std::cout << Bm;
    std::cout << "" << std::endl;

    Eigen::Matrix4i Cm = Eigen::Matrix4i::Zero();

    int Mdim = 4;
    int Ndim = 4;
    int Pdim = 4;

    const int szA = 16;
    const int szB = 16;
    const int szC = 16;
    
    // Get available platforms
    std::vector<Platform> platforms;
    cl::Platform::get(&platforms);   
    
    std::cout << "got platforms" << std::endl;
    
    // Select the default platform and create a context using this platform and the GPU
    cl_context_properties cps[3] = {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)(platforms[0])(),
        0
    };
    
    std::cout << "selected default platform" << std::endl;
    
    Context context = Context(CL_DEVICE_TYPE_GPU, cps, NULL, NULL, &err);        

    std::cout << "Context Made" << std::endl;

    // Get a list of devices on this platform
    std::vector<Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();

    if(devices.size() < 1){
        std::cout << "No GPU devices found" << std::endl;
    }
        
    std::cout << "Device Found" << std::endl;

        
    // Create a command queue and use the first device
    CommandQueue queue = CommandQueue(context, devices[0], 0, &err);
    
    std::cout << "Command Queue created" << std::endl;

    // Make program of the source code in the context
    Program program = Program(context, sources);
    if (err != CL_SUCCESS)
    {
       std::cout << "Error: Failed to create compute program!\n" << std::endl;
    }
    
    std::cout << "Program created" << std::endl;
        
    // Build program for these specific devices
    program.build(devices);
        
    std::cout << "Program Built" << std::endl;
        
    // Make kernel
    Kernel kernel(program, "iMatMult", &err);
    
    if (err != CL_SUCCESS)
    {
       std::cout << "Error: Failed to create compute kernel!\n" << std::endl;
    }
        
    std::cout << "Kernel made" << std::endl;

    // Create memory buffers
    Buffer bufferA = Buffer(context, CL_MEM_READ_ONLY, szA * sizeof(int), NULL, &err);
    Buffer bufferB = Buffer(context, CL_MEM_READ_ONLY, szB * sizeof(int), NULL, &err);
    Buffer bufferC = Buffer(context, CL_MEM_WRITE_ONLY, szC * sizeof(int), NULL, &err);
    
    std::cout << "Buffers initialized" << std::endl;

    // Copy lists A and B to the memory buffers
    queue.enqueueWriteBuffer(bufferA, CL_TRUE, 0, szA * sizeof(int), Am.data());
    queue.enqueueWriteBuffer(bufferB, CL_TRUE, 0, szB * sizeof(int), Bm.data());
    
    std::cout << "Buffers written" << std::endl;

    err = kernel.setArg(0, sizeof(int), &Mdim);
    err = kernel.setArg(1, sizeof(int), &Ndim);
    err = kernel.setArg(2, sizeof(int), &Pdim);

    err = kernel.setArg(3, bufferA);
    err = kernel.setArg(4, bufferB);
    err = kernel.setArg(5, bufferC);
        
    // Run the kernel on specific ND range
    err = queue.enqueueNDRangeKernel(kernel, NullRange, 
                                    NDRange(Mdim, Ndim), NullRange);
    std::cout << "kernel completed" << std::endl;
        
    // Read buffer C into a local list        
    err = queue.enqueueReadBuffer(bufferC, CL_TRUE, 0, 
                                szC * sizeof(int), Cm.data());
                                
    std::cout << "Completed Buffer Read" << std::endl;

    std::cout << "Result" << std::endl;
    std::cout << Cm << std::endl;
}