Another roadblock. After one iteration the neighbor count gets corrupted and starts spitting out random negative values.

Conway_OpenCL/Conway.hpp

@@ -242,7 +242,7 @@ int main(int argc, char* argv[])
 
 	// Setup the rng
 	std::mt19937 rng(time(NULL));
-	std::uniform_int_distribution<int> rgen(0, 4); // 25% chance
+	std::uniform_int_distribution<int> rgen(0, 2); // 25% chance
 
 	// Init the grids
 	unsigned char* node_grid = new unsigned char[GRID_WIDTH * GRID_HEIGHT];
@@ -262,10 +262,20 @@ int main(int argc, char* argv[])
 
 	for (int i = 0; i < GRID_WIDTH * GRID_HEIGHT * 4; i += 4) {
 
-		pixel_array[i] = i % 255; // R?
+		if (node_grid[i / 4] == 1) {
-		pixel_array[i + 1] = 70; // G?
+			pixel_array[i] = 0; // R?
-		pixel_array[i + 2] = 100; // B?
+			pixel_array[i + 1] = 0; // G?
-		pixel_array[i + 3] = 100; // A?
+			pixel_array[i + 2] = 0; // B?
+			pixel_array[i + 3] = 0; // A?	
+		}
+		else {
+			pixel_array[i] = 0; // R?
+			pixel_array[i + 1] = 0; // G?
+			pixel_array[i + 2] = 0; // B?
+			pixel_array[i + 3] = 0; // A?	
+			
+		}
+		
 	}
 
 	GLuint texture;
@@ -302,9 +312,9 @@ int main(int argc, char* argv[])
 
 	int err = 0;
 	
-	cl_mem frontWriteBuffer = clCreateFromGLTexture(context , CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, texture, &err);
+	cl_mem frontWriteBuffer = clCreateFromGLTexture(context , CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, texture, &err);
-	cl_mem frontReadBuffer = clCreateFromGLTexture(context, CL_MEM_READ_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 backBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, GRID_WIDTH * GRID_HEIGHT * sizeof(char), (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);
@@ -317,7 +327,7 @@ int main(int argc, char* argv[])
 	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, 0, sizeof(cl_mem), (void *)&frontWriteBuffer);
 	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);
@@ -336,14 +346,15 @@ int main(int argc, char* argv[])
 		// Work size, for each y line
 		size_t global_work_size[1] = { WORKER_SIZE };
 
-		status = clEnqueueAcquireGLObjects(commandQueue, 1, &frontReadBuffer, 0, 0, 0);
+		//status = clEnqueueAcquireGLObjects(commandQueue, 1, &frontReadBuffer, 0, 0, 0);
-		status = clEnqueueAcquireGLObjects(commandQueue, 1, &frontWriteBuffer, 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);
+		clFinish(commandQueue);
 		status = clEnqueueNDRangeKernel(commandQueue, align_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
 
-		status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontReadBuffer, 0, NULL, NULL);
+		//status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontReadBuffer, 0, NULL, NULL);
-		status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontWriteBuffer, 0, NULL, NULL);
+		//status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontWriteBuffer, 0, NULL, NULL);
 
 		clFinish(commandQueue);
 
@@ -366,14 +377,12 @@ int main(int argc, char* argv[])
 		// Render
 		glfwSwapBuffers(gl_window);
 
-
-
 	}
 
 	glfwTerminate();
 	
 	// Release the buffers
-	status = clReleaseMemObject(frontReadBuffer);
+//	status = clReleaseMemObject(frontReadBuffer);
 	status = clReleaseMemObject(workerCountBuffer);
 	status = clReleaseMemObject(gridWidthBuffer);
 	status = clReleaseMemObject(gridHeightBuffer);

Conway_OpenCL/conway_align.cl

@@ -1,6 +1,6 @@
 __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;
+	const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE  | CLK_FILTER_NEAREST;
 
 	// Caclulate the start and end range that this worker will be calculating
 	int data_length = *grid_width * *grid_height;
@@ -11,8 +11,11 @@ __kernel void conway_align(__read_only image2d_t front_image, __global char* bac
 
 	for (int i = start_range; i < end_range; i++){
 		
-			uint4 pixel;
+			int2 pixelcoord = (int2) (i % *grid_width, i / *grid_height);
-			pixel = read_imageui(front_image, sampler, (int2)(i,get_global_id(0)));
+			int4 pixel;
+			pixel = read_imagei(front_image, sampler, pixelcoord);
+
+			char q = pixel.w / 255;
 
 			back_image[i] = pixel.w / 255;
 	}

Conway_OpenCL/conway_compute.cl

@@ -1,8 +1,8 @@
 __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)
 {
 
-	float4 black = (float4)(.49, .68, .81, 1);
+	float4 black = (float4)(0.49, 0.68, 0.81, 1.0);
-	float4 white = (float4)(.49, .68, .71, .3);
+	float4 white = (float4)(1.0, 1.00, 1.0, 0.5);
 
 	// Caclulate the start and end range that this worker will be calculating
 	int data_length = *grid_width * *grid_height;
@@ -26,10 +26,10 @@ __kernel void conway_compute(__write_only image2d_t front_image, __global char*
 		// Top right
 		neighbors += back_image[i - *grid_width + 1];
 
-		// Right
+		/// Right
 		neighbors += back_image[i + 1];
 
-		// Bottom Right
+		/// Bottom Right
 		neighbors += back_image[i + *grid_width + 1];
 
 		// Bottom
@@ -47,13 +47,14 @@ __kernel void conway_compute(__write_only image2d_t front_image, __global char*
 		// push living status to the padded second char
 
 
-		write_imagef(front_image, pixelcoord, black);
+		//write_imagef(front_image, pixelcoord, black);
-
+		
-		if (neighbors == 3 || (neighbors == 2 && back_image[i])){
+		if (neighbors == 3 || (neighbors == 2 && back_image[i] == 1) || back_image[i] < 0){
 			write_imagef(front_image, pixelcoord, white);
 		}
 	
-		//else
+		else{
-			//write_imagei(front_image, pixelcoord, white);
+			write_imagef(front_image, pixelcoord, black);
+		}
 	}
 }