From 08fd8b2398d0a13d2fd637a711f997bec4e6c64a Mon Sep 17 00:00:00 2001 From: MitchellHansen Date: Fri, 15 Jan 2016 15:07:38 -0800 Subject: [PATCH] Revert --- Conway_OpenCL/Conway.hpp | 54 +++++++++++++-------------------- Conway_OpenCL/conway_align.cl | 17 +++-------- Conway_OpenCL/conway_compute.cl | 25 +++++++-------- 3 files changed, 37 insertions(+), 59 deletions(-) diff --git a/Conway_OpenCL/Conway.hpp b/Conway_OpenCL/Conway.hpp index 172fd93..3669fb8 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, 2); // 25% chance + std::uniform_int_distribution rgen(0, 4); // 25% chance // Init the grids unsigned char* node_grid = new unsigned char[GRID_WIDTH * GRID_HEIGHT]; @@ -262,20 +262,10 @@ int main(int argc, char* argv[]) for (int i = 0; i < GRID_WIDTH * GRID_HEIGHT * 4; i += 4) { - 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? - - } - + pixel_array[i] = i % 255; // R? + pixel_array[i + 1] = 70; // G? + pixel_array[i + 2] = 100; // B? + pixel_array[i + 3] = 100; // A? } GLuint texture; @@ -312,9 +302,9 @@ int main(int argc, char* argv[]) int err = 0; - 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 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); @@ -327,7 +317,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 *)&frontWriteBuffer); + 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); @@ -335,31 +325,27 @@ int main(int argc, char* argv[]) // ===================================== Loop ================================================================== - int i = 0; + while (!glfwWindowShouldClose(gl_window)) { - //glClearColor(0.2f, 0.3f, 0.3f, 1.0f); - //glClear(GL_COLOR_BUFFER_BIT); + glClearColor(0.2f, 0.3f, 0.3f, 1.0f); + glClear(GL_COLOR_BUFFER_BIT); // ======================================= OpenCL Shtuff =================================================== // Work size, for each y line size_t global_work_size[1] = { WORKER_SIZE }; - glFinish(); + 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); - - //clEnqueueBarrier(commandQueue); - //status = clEnqueueNDRangeKernel(commandQueue, align_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL); + 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 = clEnqueueReleaseGLObjects(commandQueue, 1, &frontReadBuffer, 0, NULL, NULL); status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontWriteBuffer, 0, NULL, NULL); - //clEnqueueBarrier(commandQueue); - - + clFinish(commandQueue); // ======================================= Rendering Shtuff ================================================= @@ -375,17 +361,19 @@ int main(int argc, char* argv[]) glDrawElements(GL_TRIANGLES, 6, GL_UNSIGNED_INT, 0); glBindVertexArray(0); - + glFinish(); // 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 90160b2..195c878 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 unsigned char* back_image, __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_NONE | CLK_FILTER_NEAREST; + 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; @@ -11,17 +11,10 @@ __kernel void conway_align(__read_only image2d_t front_image, __global unsigned for (int i = start_range; i < end_range; i++){ - int2 pixelcoord = (int2) (i % *grid_width, i / *grid_height); - int4 pixel; - pixel = read_imagei(front_image, sampler, pixelcoord); + uint4 pixel; + pixel = read_imageui(front_image, sampler, (int2)(i,get_global_id(0))); - if (pixel.w > 200){ - back_image[i] = 0; - } - else - back_image[i] = 0; - - + back_image[i] = pixel.w / 255; } } \ No newline at end of file diff --git a/Conway_OpenCL/conway_compute.cl b/Conway_OpenCL/conway_compute.cl index adf0aff..b6d2294 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 unsigned char* back_image, __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) { - float4 black = (float4)(1.0, 0.0, 0.0, 1.0); - float4 white = (float4)(0.0, 0.0, 1.0, 0.0); + float4 black = (float4)(.49, .68, .81, 1); + float4 white = (float4)(.49, .68, .71, .3); // Caclulate the start and end range that this worker will be calculating int data_length = *grid_width * *grid_height; @@ -16,8 +16,6 @@ __kernel void conway_compute(__write_only image2d_t front_image, __global unsign for (int i = start_range; i < end_range; i++){ - unsigned char im = back_image[i]; - int2 pixelcoord = (int2) (i % *grid_width, i / *grid_height); // add all 8 blocks to neghbors @@ -28,10 +26,10 @@ __kernel void conway_compute(__write_only image2d_t front_image, __global unsign // 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 @@ -49,14 +47,13 @@ __kernel void conway_compute(__write_only image2d_t front_image, __global unsign // push living status to the padded second char - //write_imagef(front_image, pixelcoord, black); - - if (neighbors == 3 || (neighbors == 2 && back_image[i] == 1) || (i % 10) == 1){ - write_imagef(front_image, pixelcoord, black); - } - - else{ + 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); } } \ No newline at end of file