diff --git a/Conway_OpenCL/Conway.hpp b/Conway_OpenCL/Conway.hpp index 3669fb8..0146f07 100644 --- a/Conway_OpenCL/Conway.hpp +++ b/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 rgen(0, 4); // 25% chance + std::uniform_int_distribution 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? - pixel_array[i + 1] = 70; // G? - pixel_array[i + 2] = 100; // B? - pixel_array[i + 3] = 100; // A? + if (node_grid[i / 4] == 1) { + pixel_array[i] = 0; // R? + pixel_array[i + 1] = 0; // G? + 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 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 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 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, &frontWriteBuffer, 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); + 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, &frontWriteBuffer, 0, NULL, NULL); + //status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontReadBuffer, 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); diff --git a/Conway_OpenCL/conway_align.cl b/Conway_OpenCL/conway_align.cl index 195c878..7faab6d 100644 --- a/Conway_OpenCL/conway_align.cl +++ b/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; - pixel = read_imageui(front_image, sampler, (int2)(i,get_global_id(0))); + int2 pixelcoord = (int2) (i % *grid_width, i / *grid_height); + int4 pixel; + pixel = read_imagei(front_image, sampler, pixelcoord); + + char q = pixel.w / 255; back_image[i] = pixel.w / 255; } diff --git a/Conway_OpenCL/conway_compute.cl b/Conway_OpenCL/conway_compute.cl index b6d2294..d1bf5d6 100644 --- a/Conway_OpenCL/conway_compute.cl +++ b/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 white = (float4)(.49, .68, .71, .3); + float4 black = (float4)(0.49, 0.68, 0.81, 1.0); + 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); - - if (neighbors == 3 || (neighbors == 2 && back_image[i])){ + //write_imagef(front_image, pixelcoord, black); + + if (neighbors == 3 || (neighbors == 2 && back_image[i] == 1) || back_image[i] < 0){ write_imagef(front_image, pixelcoord, white); } - //else - //write_imagei(front_image, pixelcoord, white); + else{ + write_imagef(front_image, pixelcoord, black); + } } } \ No newline at end of file