It draws!!

Conway_OpenCL/Conway.hpp

@@ -74,7 +74,7 @@ int main(int argc, char* argv[])
 	int WINDOW_Y = 1000;
 	int GRID_WIDTH = WINDOW_X;
 	int GRID_HEIGHT = WINDOW_Y;
-	int WORKER_SIZE = 2000;
+	int WORKER_SIZE = 1000;
 
 	// ======================================= Setup OpenGL =======================================================
 
@@ -97,7 +97,8 @@ int main(int argc, char* argv[])
 
 	glClearColor(0.2f, 0.3f, 0.3f, 1.0f);
 
-	Shader ourShader("Z:\\VS_Projects\\Conway_OpenCL\\Conway_OpenCL\\vertex_shader.sh", "Z:\\VS_Projects\\Conway_OpenCL\\Conway_OpenCL\\fragment_shader.sh");
+	// Stolen from LearnOpenGL, aint got no time for that
+	Shader vert_frag_shaders("Z:\\VS_Projects\\Conway_OpenCL\\Conway_OpenCL\\vertex_shader.sh", "Z:\\VS_Projects\\Conway_OpenCL\\Conway_OpenCL\\fragment_shader.sh");
 
 	GLfloat vertices[] = {
 		// Positions          // Colors           // Texture Coords
@@ -233,7 +234,7 @@ int main(int argc, char* argv[])
 
 	// Now create the kernels
 	cl_kernel compute_kernel = clCreateKernel(compute_program, "conway_compute", NULL);
-	cl_kernel back_kernel = clCreateKernel(align_program, "conway_align", NULL);
+	cl_kernel align_kernel = clCreateKernel(align_program, "conway_align", NULL);
 
 
 
@@ -244,14 +245,14 @@ int main(int argc, char* argv[])
 	std::uniform_int_distribution<int> rgen(0, 4); // 25% chance
 
 	// Init the grids
-	unsigned char* front_grid = new unsigned char[GRID_WIDTH * GRID_HEIGHT];
+	unsigned char* node_grid = new unsigned char[GRID_WIDTH * GRID_HEIGHT];
 
 	for (int i = 0; i < GRID_WIDTH * GRID_HEIGHT; i++) {
 		if (rgen(rng) == 1) {
-			front_grid[i] = 1;
+			node_grid[i] = 1;
 		}
 		else {
-			front_grid[i] = 0;
+			node_grid[i] = 0;
 		}
 	}
 
@@ -293,7 +294,7 @@ int main(int argc, char* argv[])
 	glGenerateMipmap(GL_TEXTURE_2D);
 
 
-	//delete pixel_array;
+	delete pixel_array;
 
 
 
@@ -301,58 +302,48 @@ int main(int argc, char* argv[])
 
 	int err = 0;
 	
-	cl_mem frontBuffer = clCreateFromGLTexture(context , CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, texture, &err);
+	cl_mem frontWriteBuffer = clCreateFromGLTexture(context , CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, texture, &err);
+	cl_mem frontReadBuffer = clCreateFromGLTexture(context, CL_MEM_READ_ONLY, GL_TEXTURE_2D, 0, texture, &err);
+	cl_mem backBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int), (void*)node_grid, &err);
 
 	cl_mem workerCountBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int), &WORKER_SIZE, &err);
 	cl_mem gridWidthBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int), &GRID_WIDTH, &err);
 	cl_mem gridHeightBuffer = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(int), &GRID_HEIGHT, &err);
 
-	// Kernel args
+	// Compute kernel args
-	status = clSetKernelArg(compute_kernel, 0, sizeof(cl_mem), (void *)&frontBuffer);
+	status = clSetKernelArg(compute_kernel, 0, sizeof(cl_mem), (void *)&frontWriteBuffer);
-	status = clSetKernelArg(compute_kernel, 1, sizeof(cl_mem), (void *)&workerCountBuffer);
+	status = clSetKernelArg(compute_kernel, 1, sizeof(cl_mem), (void *)&backBuffer);
-	status = clSetKernelArg(compute_kernel, 2, sizeof(cl_mem), (void *)&gridWidthBuffer);
+	status = clSetKernelArg(compute_kernel, 2, sizeof(cl_mem), (void *)&workerCountBuffer);
-	status = clSetKernelArg(compute_kernel, 3, sizeof(cl_mem), (void *)&gridHeightBuffer);
+	status = clSetKernelArg(compute_kernel, 3, sizeof(cl_mem), (void *)&gridWidthBuffer);
+	status = clSetKernelArg(compute_kernel, 4, sizeof(cl_mem), (void *)&gridHeightBuffer);
 
+	status = clSetKernelArg(align_kernel, 0, sizeof(cl_mem), (void *)&frontReadBuffer);
+	status = clSetKernelArg(align_kernel, 1, sizeof(cl_mem), (void *)&backBuffer);
+	status = clSetKernelArg(align_kernel, 2, sizeof(cl_mem), (void *)&workerCountBuffer);
+	status = clSetKernelArg(align_kernel, 3, sizeof(cl_mem), (void *)&gridWidthBuffer);
+	status = clSetKernelArg(align_kernel, 4, sizeof(cl_mem), (void *)&gridHeightBuffer);
 
 
 	// ===================================== Loop ==================================================================
 
 	while (!glfwWindowShouldClose(gl_window)) {
 
-		// Clear the colorbuffer
 		glClearColor(0.2f, 0.3f, 0.3f, 1.0f);
 		glClear(GL_COLOR_BUFFER_BIT);
 
-		//glfwPollEvents();
-		//glClear(GL_COLOR_BUFFER_BIT);
-
 		// ======================================= OpenCL Shtuff ===================================================
 
-		// Update the data in GPU memory
-		//status = clEnqueueWriteBuffer(commandQueue, frontBuffer, CL_TRUE, 0, GRID_WIDTH * GRID_HEIGHT * 2 * sizeof(char), (void*)grid, NULL, 0, NULL);
-
 		// Work size, for each y line
-		size_t global_work_size[1] = { 10 };
+		size_t global_work_size[1] = { WORKER_SIZE };
-
-		glActiveTexture(GL_TEXTURE0);
-		glBindTexture(GL_TEXTURE_2D, texture);
-		glUniform1i(glGetUniformLocation(ourShader.Program, "ourTexture1"), 0);
-
-		// Draw the triangle
-		ourShader.Use();
-		glBindVertexArray(VAO);
-		glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0);
-		glBindVertexArray(0);
-
-		glFinish();
 
-		status = clEnqueueAcquireGLObjects(commandQueue, 1, &frontBuffer, 0, 0, 0);
+		status = clEnqueueAcquireGLObjects(commandQueue, 1, &frontReadBuffer, 0, 0, 0);
+		status = clEnqueueAcquireGLObjects(commandQueue, 1, &frontWriteBuffer, 0, 0, 0);
 
 		status = clEnqueueNDRangeKernel(commandQueue, compute_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
+		status = clEnqueueNDRangeKernel(commandQueue, align_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
 
-		//status = clEnqueueReadBuffer(commandQueue, frontBuffer, CL_TRUE, 0, GRID_WIDTH * GRID_HEIGHT * 4 * sizeof(unsigned char), (void*)pixel_array, 0, NULL, NULL);
+		status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontReadBuffer, 0, NULL, NULL);
-
+		status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontWriteBuffer, 0, NULL, NULL);
-		status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontBuffer, 0, NULL, NULL);
 
 		clFinish(commandQueue);
 
@@ -360,11 +351,19 @@ int main(int argc, char* argv[])
 
 		glfwPollEvents();
 
-		// Render
+		glActiveTexture(GL_TEXTURE0);
+		glBindTexture(GL_TEXTURE_2D, texture);
+		glUniform1i(glGetUniformLocation(vert_frag_shaders.Program, "pixel_texture"), 0);
 
+		// Draw the triangle
+		vert_frag_shaders.Use();
+		glBindVertexArray(VAO);
+		glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0);
+		glBindVertexArray(0);
 
+		glFinish();
 
-		// Swap the screen buffers
+		// Render
 		glfwSwapBuffers(gl_window);
 
 
@@ -374,7 +373,7 @@ int main(int argc, char* argv[])
 	glfwTerminate();
 	
 	// Release the buffers
-	status = clReleaseMemObject(frontBuffer);
+	status = clReleaseMemObject(frontReadBuffer);
 	status = clReleaseMemObject(workerCountBuffer);
 	status = clReleaseMemObject(gridWidthBuffer);
 	status = clReleaseMemObject(gridHeightBuffer);

Conway_OpenCL/conway_align.cl

@@ -1,7 +1,20 @@
-__kernel void conway_align(__global unsigned char* front_grid, __global unsigned char* rear_grid, __global unsigned char* pixel_out, __global int* num_workers, __global int* grid_width, __global int* grid_height){
+__kernel void conway_align(__read_only image2d_t front_image, __global char* back_image, __global int* num_workers, __global int* grid_width, __global int *grid_height)
+{
+	const sampler_t sampler=CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
 
+	// Caclulate the start and end range that this worker will be calculating
+	int data_length = *grid_width * *grid_height;
 
+	int start_range = (data_length / *num_workers) * get_global_id(0);
+	int end_range = (data_length / *num_workers) * (get_global_id(0) + 1);
 
 
+	for (int i = start_range; i < end_range; i++){
+		
+			uint4 pixel;
+			pixel = read_imageui(front_image, sampler, (int2)(i,get_global_id(0)));
+
+			back_image[i] = pixel.w / 255;
+	}
 
 }

Conway_OpenCL/conway_compute.cl

@@ -1,21 +1,59 @@
-__kernel void conway_compute(__write_only image2d_t front_grid, __global int* num_workers, __global int* grid_width, __global int *grid_height)
+__kernel void conway_compute(__write_only image2d_t front_image, __global char* back_image, __global int* num_workers, __global int* grid_width, __global int *grid_height)
 {
 
-//int width = *grid_width;
+	float4 black = (float4)(.49, .68, .81, 1);
-//int height = grid_height;
+	float4 white = (float4)(.49, .68, .71, .3);
 
-	for (int i = 0; i < 90000; i ++){
+	// Caclulate the start and end range that this worker will be calculating
-	int2 pixelcoord = (int2) (i % *grid_width, i / *grid_height);
+	int data_length = *grid_width * *grid_height;
-	//if (pixelcoord.x < width && pixelcoord.y < height)
-	//{   
-		//float4 pixel = read_imagef(image1, sampler, (int2)(pixelcoord.x, pixelcoord.y));
-    int4 black = (int4)(0,0,0,0);
 
+	int start_range = (data_length / *num_workers) * get_global_id(0);
+	int end_range = (data_length / *num_workers) * (get_global_id(0) + 1);
 	
-    //write_imagef(front_grid, pixelcoord, black);
+	// x, y + 1
 
-	write_imagei(front_grid, pixelcoord, black);
+	int neighbors = 0;
 
-}
+	for (int i = start_range; i < end_range; i++){
-//}
+
+		int2 pixelcoord = (int2) (i % *grid_width, i / *grid_height);
+
+		// add all 8 blocks to neghbors
+		neighbors = 0;
+		// Top
+		neighbors += back_image[i - *grid_width];
+
+		// Top right
+		neighbors += back_image[i - *grid_width + 1];
+
+		// Right
+		neighbors += back_image[i + 1];
+
+		// Bottom Right
+		neighbors += back_image[i + *grid_width + 1];
+
+		// Bottom
+		neighbors += back_image[i + *grid_width];
+
+		// Bottom Left
+		neighbors += back_image[i + *grid_width - 1];
+
+		// Left
+		neighbors += back_image[i - 1];
+
+		// Top left
+		neighbors += back_image[i - *grid_width - 1];
+
+		// push living status to the padded second char
+
+
+		write_imagef(front_image, pixelcoord, black);
+
+		if (neighbors == 3 || (neighbors == 2 && back_image[i])){
+			write_imagef(front_image, pixelcoord, white);
+		}
+	
+		//else
+			//write_imagei(front_image, pixelcoord, white);
+	}
 }

Conway_OpenCL/fragment_shader.sh

@@ -5,11 +5,11 @@ in vec2 TexCoord;
 out vec4 color;
 
 // Texture samplers
-uniform sampler2D ourTexture1;
+uniform sampler2D pixel_texture;
 
 void main()
 {
 	// Linearly interpolate between both textures (second texture is only slightly combined)
 	//color = vec4(1.0f, 0.5f, 0.2f, 1.0f);
-	color = texture(ourTexture1, TexCoord);
+	color = texture(pixel_texture, TexCoord);
 }