// Conway's Game of Life
//
// rules:
// 1. Every cell is labeled "alive" or "dead". Each updates itself,
//    synchronously, on each iteration.
// 2. Dead cells with exactly 3 live neighbors (including diagonals)
//    come to life.
// 3. Live cells with fewer than 2 live neighbors die of loneliness.
// 4. Live cells with more than 3 live neighbors die of overcrowding.
//
// Command line argument is the probability that the cell is initially alive.
//

#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <oclUtils.h>
#include <GL/glew.h>
#include <GL/glut.h>
#include <GL/glx.h>
#include <CL/cl_gl.h>

#define FULL_INTEROP 0

#define WIDTH 768
#define HEIGHT 768
#define DATA_SIZE (WIDTH*HEIGHT*4*sizeof(GLubyte)) 

#define OGL_PBO 1

// OpenCL vars
cl_platform_id myplatform;
cl_context mycontext;
cl_device_id *mydevice;
cl_command_queue mycq;
cl_kernel mykernel;
cl_program myprogram;
cl_image_format myimformat;
cl_mem oclpbo[2];
cl_int err;
char* oclsource; 
size_t gwsize[2] = {WIDTH,HEIGHT}; 
size_t lwsize[2] = {8,8}; 

unsigned char *host_image;
unsigned char *red_image;
int bindex = 1;
int tid[] = {1,2};

void do_kernel()
{
cl_event wlist[1];
// flip this (current buffer index) on each call
bindex = 1-bindex;
clSetKernelArg(mykernel,0,sizeof(cl_mem),(void *)&oclpbo[bindex]);
clSetKernelArg(mykernel,1,sizeof(cl_mem),(void *)&oclpbo[1-bindex]);

clEnqueueNDRangeKernel(mycq,mykernel,2,NULL,gwsize,lwsize,0,0,&wlist[0]);
clWaitForEvents(1,wlist);
}

void mydisplayfunc()
{
const size_t start[3] = {0,0,0};
const size_t scope[3] = {WIDTH,HEIGHT,1};
void *ptr;
int err;

#if FULL_INTEROP
glFinish();
err=clEnqueueAcquireGLObjects(mycq,2,&oclpbo[0],0,0,0);
#endif

do_kernel();

#if FULL_INTEROP
err=clEnqueueReleaseGLObjects(mycq,2,&oclpbo[0],0,0,0);
clFinish(mycq);

glClear(GL_COLOR_BUFFER_BIT);
glColor3f(1.0,0.0,0.0);
glEnable(GL_TEXTURE_2D);
glBindTexture(GL_TEXTURE_2D,tid[1-bindex]);
glBegin(GL_QUADS);
// Just draw a square with this texture on it.
glTexCoord2f(0.0,1.0);
glVertex2f(-1.0,-1.0);
glTexCoord2f(1.0,1.0);
glVertex2f(1.0,-1.0);
glTexCoord2f(1.0,0.0);
glVertex2f(1.0,1.0);
glTexCoord2f(0.0,0.0);
glVertex2f(-1.0,1.0);
glEnd();
glDisable(GL_TEXTURE_2D);
glBindTexture(GL_TEXTURE_2D,0);
glFlush();
#else
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, OGL_PBO);
// map the current pbo to user's address space at "ptr"
ptr = glMapBuffer(GL_PIXEL_UNPACK_BUFFER, GL_READ_WRITE);
// use blocking call here
clEnqueueReadImage(mycq,oclpbo[1-bindex],CL_TRUE,start,scope,0,0,ptr,0,NULL,NULL);
glUnmapBuffer(GL_PIXEL_UNPACK_BUFFER); 
glClear(GL_COLOR_BUFFER_BIT);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, OGL_PBO);
glDrawPixels(WIDTH,HEIGHT,GL_RGBA,GL_UNSIGNED_BYTE,0);
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0);
#endif
glutSwapBuffers();
usleep(50000);
glutPostRedisplay();
}

void InitGL(int argc, char** argv)
{
glutInit(&argc,argv);
glutInitDisplayMode(GLUT_RGBA|GLUT_DOUBLE);
glutInitWindowSize(WIDTH,HEIGHT);
glutInitWindowPosition(100,50);
glutCreateWindow("Life");
glDisable(GL_DEPTH_TEST);
glClearColor(0.1,0.2,0.35,1.0);
glewInit();
#if FULL_INTEROP
glBindTexture(GL_TEXTURE_2D,1);
glTexImage2D(GL_TEXTURE_2D,0,GL_RGBA8,WIDTH,HEIGHT,0,GL_RGBA,
        GL_UNSIGNED_BYTE,host_image);
free(host_image);
glTexParameterf(GL_TEXTURE_2D,GL_TEXTURE_MIN_FILTER,GL_NEAREST);
glTexParameterf(GL_TEXTURE_2D,GL_TEXTURE_MAG_FILTER,GL_NEAREST);
glTexEnvf(GL_TEXTURE_ENV,GL_TEXTURE_ENV_MODE,GL_REPLACE);
glBindTexture(GL_TEXTURE_2D,0);
glBindTexture(GL_TEXTURE_2D,2);
glTexImage2D(GL_TEXTURE_2D,0,GL_RGBA8,WIDTH,HEIGHT,0,GL_RGBA,
        GL_UNSIGNED_BYTE,red_image);
free(red_image);
glTexParameterf(GL_TEXTURE_2D,GL_TEXTURE_MIN_FILTER,GL_NEAREST);
glTexParameterf(GL_TEXTURE_2D,GL_TEXTURE_MAG_FILTER,GL_NEAREST);
glTexEnvf(GL_TEXTURE_ENV,GL_TEXTURE_ENV_MODE,GL_REPLACE);
glBindTexture(GL_TEXTURE_2D,0);
glFlush();
#endif

return;
}

double genrand()
{
return(((double)(random())+1.0)/((double)(RAND_MAX)+2.0));
}

void init_host_image(float frac_on)
{
int i;

host_image = (unsigned char *)calloc(1,DATA_SIZE);
red_image = (unsigned char *)calloc(1,DATA_SIZE);
srandom(123456789);

/*  Create an initial texture.  White is "alive"; cyan is "dead". */
for(i=0;i<WIDTH*HEIGHT;i++){
        if(genrand()<frac_on){
                /* white */
                host_image[4*i+0] = 255;
                host_image[4*i+1] = 255;
                host_image[4*i+2] = 255;
                host_image[4*i+3] = 0;
                }
        else{
                /* dark cyan */
                host_image[4*i+0] = 0;
                host_image[4*i+1] = 128;
                host_image[4*i+2] = 128;
                host_image[4*i+3] = 0;
                }
	/* red is test */
	red_image[4*i+0] = 255;
       	red_image[4*i+1] = 0;
       	red_image[4*i+2] = 0;
       	red_image[4*i+3] = 0;
        }
}

void cleanup()
{
clReleaseKernel(mykernel);
clReleaseProgram(myprogram);
clReleaseCommandQueue(mycq);
clReleaseMemObject(oclpbo[0]);
clReleaseMemObject(oclpbo[1]);
clReleaseContext(mycontext);
exit(0);
}

void getout(unsigned char key, int x, int y)
{
switch(key) {
        case 'q':               // escape quits
                cleanup();
        default:
                break;
    }
}

int main(int argc,char **argv)
{
size_t program_length;
unsigned int width, height, gpudevcount;
int err;

// load ppm image of initial population with padding inserted during load
init_host_image(atof(argv[1]));

InitGL(argc, argv); 

err = oclGetPlatformID(&myplatform);
err = clGetDeviceIDs(myplatform,CL_DEVICE_TYPE_GPU,0,NULL,&gpudevcount);
fprintf(stderr,"GPU device count %d\n",gpudevcount);

mydevice = new cl_device_id[gpudevcount];
err = clGetDeviceIDs(myplatform,CL_DEVICE_TYPE_GPU,gpudevcount,mydevice,NULL);

cl_context_properties props[] = {
        CL_GL_CONTEXT_KHR, (cl_context_properties)glXGetCurrentContext(),
        CL_GLX_DISPLAY_KHR, (cl_context_properties)glXGetCurrentDisplay(),
        CL_CONTEXT_PLATFORM, (cl_context_properties)myplatform,
        0};
mycontext = clCreateContext(props,1,&mydevice[0],NULL,NULL,&err);
fprintf(stderr,"context creation %d\n",err);

mycq = clCreateCommandQueue(mycontext,mydevice[0],0,&err);
fprintf(stderr,"command queue %d\n",err);

oclsource = oclLoadProgSource("newlife.cl", "", &program_length);
myprogram = clCreateProgramWithSource(mycontext,1,(const char **)&oclsource,
	&program_length, &err);
clBuildProgram(myprogram, 0, NULL, NULL, NULL, NULL);
mykernel = clCreateKernel(myprogram, "flip", &err);
fprintf(stderr,"err is %d\n",err);

#if FULL_INTEROP
oclpbo[0] = clCreateFromGLTexture2D(mycontext,CL_MEM_READ_WRITE,GL_TEXTURE_2D,0,1,&err);
oclpbo[1] = clCreateFromGLTexture2D(mycontext,CL_MEM_READ_WRITE,GL_TEXTURE_2D,0,2,&err);
#else
glBindBuffer(GL_PIXEL_UNPACK_BUFFER, OGL_PBO);
glBufferData(GL_PIXEL_UNPACK_BUFFER, DATA_SIZE, NULL, GL_DYNAMIC_DRAW);

myimformat.image_channel_order = CL_RGBA;
myimformat.image_channel_data_type = CL_UNSIGNED_INT8;

// pbo 0 gets the initial load from host_image
oclpbo[0] = clCreateImage2D(mycontext,CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
	&myimformat,WIDTH,HEIGHT,0,host_image,&err);
oclpbo[1] = clCreateImage2D(mycontext,CL_MEM_READ_WRITE,
	&myimformat,WIDTH,HEIGHT,0,NULL,&err);
#endif

glutDisplayFunc(mydisplayfunc);
glutKeyboardFunc(getout);
glutMainLoop();
}