#include <stdlib.h>
#include <SDL/SDL.h>
#include <GL/gl.h>
#include <GL/glu.h>
#include <GL/glext.h>
#include <GL/glx.h>
#include <CL/cl.h>
#include <CL/cl_gl.h>
#include <math.h>
#include <sys/time.h>
#include <iostream>

using namespace std;

GLXContext gl_context;
Display *display;

GLuint vbo;
GLuint indicies;

cl_context cl_c;
cl_command_queue queue;
cl_program program;
cl_kernel kernel;
cl_mem cl_m;
cl_mem cl_m2;
timeval now,old,res;

const int grid_size = 2048;

const char source[] = "#include \"kernel.cl\"";
const char *kernel_source[] = { source };

float rott = 0;

const int w = 1024;
const int h = 768;

void InitGL()
{
    glClearColor(0.0f, 0.5f, 0.0f, 0.0f);
    glViewport(0,0,w,h);
    glClearDepth(1.0f);
    glMatrixMode(GL_PROJECTION);
    glLoadIdentity();
    //glFrustum(-0.5, 500, -0.5, 500.0, 1.0, 100.0);
    gluPerspective(40.0f, (float)w/h, 0.1f, 100.0f);
    glMatrixMode(GL_MODELVIEW);
    glLoadIdentity();

    glGenBuffers(1, &vbo);
    glBindBuffer(GL_ARRAY_BUFFER, vbo);
    float *grid = new float[grid_size*grid_size*4];
    for(int i=0;i<grid_size;i++)
    {
        for(int o=0;o<grid_size;o++)
        {
            grid[i*grid_size*4+o*4] = (float)o/grid_size*10.0-5;
            grid[i*grid_size*4+o*4+1] = (float)i/grid_size*10.0-5;
            grid[i*grid_size*4+o*4+2] = sin((float)o/grid_size*6.283185308*2)+cos((float)i/grid_size*6.283185308*2);
            grid[i*grid_size*4+o*4+3] = 1.0f;
        }
    }
    glBufferData(GL_ARRAY_BUFFER, sizeof(float)*grid_size*grid_size*4, grid, GL_STREAM_DRAW);
    delete [] grid;

    unsigned int *ind = new unsigned int[(grid_size-1)*(grid_size-1)*6];
    for(int i=0;i<(grid_size-1);i++)
    {
        for(int o=0;o<(grid_size-1);o++)
        {
            ind[i*(grid_size-1)*6+o*6] = i*(grid_size)+o;
            ind[i*(grid_size-1)*6+o*6+1] = i*(grid_size)+o+1;
            ind[i*(grid_size-1)*6+o*6+2] = (i+1)*(grid_size)+o+1;
            ind[i*(grid_size-1)*6+o*6+3] = i*(grid_size)+o;
            ind[i*(grid_size-1)*6+o*6+4] = (i+1)*(grid_size)+o+1;
            ind[i*(grid_size-1)*6+o*6+5] = (i+1)*(grid_size)+o;
        }
    }
    glGenBuffers(1, &indicies);
    glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, indicies);
    glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(int)*(grid_size-1)*(grid_size-1)*6, ind, GL_STATIC_DRAW);
    delete [] ind;
    //glPolygonMode(GL_BACK, GL_LINE);

    printf("GL inicializated %d\n", glGetError());
}

void InitCL()
{
    cl_int err_code;
    cl_platform_id platform;
    err_code = clGetPlatformIDs(1, &platform, NULL);
    if(err_code)
    {
        printf("Unable get platform %d\n", err_code);
        exit(1);
    }
    cl_context_properties cl_properties[] = {CL_GL_CONTEXT_KHR, (cl_context_properties)gl_context,
        CL_GLX_DISPLAY_KHR, (cl_context_properties)display,
        CL_CONTEXT_PLATFORM, (cl_context_properties)platform, 0};
    cl_c = clCreateContextFromType(cl_properties, CL_DEVICE_TYPE_GPU, NULL, NULL, &err_code);
    if(err_code)
    {
        printf("Unable create context %d\n", err_code);
        exit(1);
    }

    cl_device_id device;
    clGetContextInfo(cl_c, CL_CONTEXT_DEVICES, sizeof(cl_device_id), &device, NULL);
    queue = clCreateCommandQueue(cl_c, device, NULL, &err_code);
    if(err_code)
    {
        printf("Unable create command queue %d\n", err_code);
        exit(1);
    }
    size_t source_len[] = { sizeof(source)-1 };
    //const char s[] = {kernel_source , 0};
    program = clCreateProgramWithSource(cl_c, 1, kernel_source, source_len, &err_code);
    if(err_code)
    {
        printf("Unable create program %d\n", err_code);
        exit(1);
    }
    err_code = clBuildProgram(program, 1, &device, "-I.", NULL, NULL);
    if(err_code)
    {
        printf("Unable build program %d\n", err_code);
    }
        char log[10000];
        clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 10000, log, NULL);
        printf("BUILD LOG:\n%s", log);
        if(err_code)exit(0);
    kernel = clCreateKernel(program, "Wave", &err_code);
    if(err_code)
    {
        printf("Unable create kernel %d\n", err_code);
    }
    size_t size;
    clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &size, NULL);
    cout << "kernel work group size: "  <<  size << endl;
    clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_COMPILE_WORK_GROUP_SIZE, sizeof(size_t), &size, NULL);
    cout << "kernel compile work group size: " << size << endl;

    cl_m2 = clCreateBuffer(cl_c, CL_MEM_READ_WRITE, sizeof(float)*grid_size*grid_size*4, NULL, &err_code);
    if(err_code)
    {
        printf("Unable create buffer %d\n", err_code);
        exit(1);
    }

}

void RunCL()
{
    cl_int err_code;
    glFinish();
    err_code = clEnqueueAcquireGLObjects(queue, 1, &cl_m, 0, NULL, NULL);
    if(err_code)printf("Unable acquire object %d\n", err_code);
    err_code = clSetKernelArg(kernel, 0, sizeof(float), &rott);
    err_code |= clSetKernelArg(kernel, 1, sizeof(int), &grid_size);
    err_code |= clSetKernelArg(kernel, 2, sizeof(cl_mem), &cl_m);
    if(err_code)printf("Unable set arguments %d\n", err_code);

    size_t work[] = {grid_size, grid_size};
    clFinish(queue);
    err_code = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, work, NULL, 0, NULL, NULL);
    //clFinish(queue);

    if(err_code)
    {
        printf("Unable run a kernel %d\n", err_code);
        exit(1);
    }
    err_code = clEnqueueReleaseGLObjects(queue, 1, &cl_m, 0, NULL, NULL);
    if(err_code)printf("Unable release object %d\n", err_code);
    clFinish(queue);
    gettimeofday(&now,NULL);
    timersub(&now, &old,&res);
    old = now;
    //printf("Cas %f\n",res.tv_sec + res.tv_usec/1000000.0f);;
}

void DrawGL()
{
    static int last_time = SDL_GetTicks();
    static int fps = 0;
    if(SDL_GetTicks() - last_time > 1000)
    {
        cout << "FPS: " << fps << endl;
        fps = 0;
        last_time = SDL_GetTicks();
    }
    fps++;
    glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
    glLoadIdentity();
    glTranslatef( 0.0f, 0.0f, -20.0f );
    glRotatef(rott, 1.0f, 0.0f, 0.0f);
    rott += 0.4f;

    /*glBegin(GL_QUADS);
    glVertex3f(0.0f, 0.0f, 0.0f);
    glVertex3f(5.0f, 0.0f, 0.0f);
    glVertex3f(5.0f, 5.0f, 0.0f);
    glVertex3f(0.0f, 5.0f, 0.0f);
    glEnd();*/

    glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, indicies);
    glBindBuffer(GL_ARRAY_BUFFER, vbo);
    glEnableClientState(GL_VERTEX_ARRAY);
    glVertexPointer(4, GL_FLOAT, 0, 0);
    glDrawElements(GL_TRIANGLES, (grid_size-1)*(grid_size-1)*6, GL_UNSIGNED_INT, NULL);
    glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, 0);
    glBindBuffer(GL_ARRAY_BUFFER, 0);
    glDisableClientState(GL_VERTEX_ARRAY);
    //printf("%d\n", glGetError());
    SDL_GL_SwapBuffers();
    //glFinish();
}

int main ( int argc, char** argv )
{
    // initialize SDL video
    if ( SDL_Init( SDL_INIT_VIDEO ) < 0 )
    {
        printf( "Unable to init SDL: %s\n", SDL_GetError() );
        return 1;
    }

    // make sure SDL cleans up before exit
    atexit(SDL_Quit);

    // create a new window
    //SDL_GL_SetAttribute(SDL_GL_MULTISAMPLEBUFFERS, 1);
    //SDL_GL_SetAttribute(SDL_GL_MULTISAMPLESAMPLES, 8);
    SDL_Surface* screen = SDL_SetVideoMode(w, h, 32, SDL_OPENGL|SDL_GL_DOUBLEBUFFER|SDL_HWSURFACE);
    if ( !screen )
    {
        printf("Unable to set 640x480 video: %s\n", SDL_GetError());
        return 1;
    }
    SDL_WM_SetCaption("OpenCL - OpenGL interoperability example", NULL);

    GLint tex_size;
    glGetIntegerv(GL_MAX_TEXTURE_SIZE, &tex_size);
    cout << "Maximum texture size " << tex_size << endl;

    gl_context = glXGetCurrentContext();
    display = glXGetCurrentDisplay();

    InitCL();
    InitGL();

    cl_int err_code;
    cl_m = clCreateFromGLBuffer(cl_c, CL_MEM_READ_WRITE, vbo, &err_code);
    size_t size;
    clGetMemObjectInfo(cl_m, CL_MEM_SIZE, sizeof(size_t), &size, NULL);

    cout << cl_m << " " << size/1024/1024 << " MiB" << endl;
    if(err_code)
    {
        printf("Unable share GL object %d\n", err_code);
    }

    // program main loop
    bool done = false;
    while (!done)
    {
        // message processing loop
        SDL_Event event;
        while (SDL_PollEvent(&event))
        {
            // check for messages
            switch (event.type)
            {
                // exit if the window is closed
            case SDL_QUIT:
                done = true;
                cout << "quit" << endl;
                break;

                // check for keypresses
            case SDL_KEYDOWN:
                {
                    if(event.key.keysym.sym == SDLK_r)
                    {
                        RunCL();
                    }
                    // exit if ESCAPE is pressed
                    if (event.key.keysym.sym == SDLK_ESCAPE)
                        done = true;
                    break;
                }
            } // end switch
        } // end of message processing

        DrawGL();
        //RunCL();
        //SDL_Delay(10);
    } // end main loop

    return 0;
}

/*********************************
kernel.cl
*********************************/
//#pragma OPENCL EXTENSION cl_khr_gl_sharing : enable
__kernel void Wave(float time, int grid_size, __global float *data)
{

    __local int ll[128];
    int i = get_global_id(0);
    int o = get_global_id(1);
    data[o*grid_size*4+i*4+2] = sin((float)o/grid_size*6.283185308f*2+time/4)+cos((float)i/grid_size*6.283185308f*2+time/4);;
	//hej[0] =         data[o*grid_size*4+i*4+2] ;
}