From 1b379e095e6c5863bdbdebec10aca90d32c1c4f7 Mon Sep 17 00:00:00 2001 From: MitchellHansen Date: Sun, 10 Jan 2016 14:22:23 -0800 Subject: [PATCH] It draws!! --- Conway_OpenCL/Conway.hpp | 79 ++++++++++++++++---------------- Conway_OpenCL/conway_align.cl | 15 +++++- Conway_OpenCL/conway_compute.cl | 64 ++++++++++++++++++++------ Conway_OpenCL/fragment_shader.sh | 4 +- 4 files changed, 106 insertions(+), 56 deletions(-) diff --git a/Conway_OpenCL/Conway.hpp b/Conway_OpenCL/Conway.hpp index 5768bde..3669fb8 100644 --- a/Conway_OpenCL/Conway.hpp +++ b/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 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 - status = clSetKernelArg(compute_kernel, 0, sizeof(cl_mem), (void *)&frontBuffer); - status = clSetKernelArg(compute_kernel, 1, sizeof(cl_mem), (void *)&workerCountBuffer); - status = clSetKernelArg(compute_kernel, 2, sizeof(cl_mem), (void *)&gridWidthBuffer); - status = clSetKernelArg(compute_kernel, 3, sizeof(cl_mem), (void *)&gridHeightBuffer); + // Compute kernel args + status = clSetKernelArg(compute_kernel, 0, sizeof(cl_mem), (void *)&frontWriteBuffer); + status = clSetKernelArg(compute_kernel, 1, sizeof(cl_mem), (void *)&backBuffer); + status = clSetKernelArg(compute_kernel, 2, sizeof(cl_mem), (void *)&workerCountBuffer); + 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 }; - - 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(); + size_t global_work_size[1] = { WORKER_SIZE }; - 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, &frontBuffer, 0, NULL, NULL); + status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontReadBuffer, 0, NULL, NULL); + status = clEnqueueReleaseGLObjects(commandQueue, 1, &frontWriteBuffer, 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); diff --git a/Conway_OpenCL/conway_align.cl b/Conway_OpenCL/conway_align.cl index 94de3c5..195c878 100644 --- a/Conway_OpenCL/conway_align.cl +++ b/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; + } } \ No newline at end of file diff --git a/Conway_OpenCL/conway_compute.cl b/Conway_OpenCL/conway_compute.cl index acbdf9b..b6d2294 100644 --- a/Conway_OpenCL/conway_compute.cl +++ b/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; -//int height = grid_height; + float4 black = (float4)(.49, .68, .81, 1); + float4 white = (float4)(.49, .68, .71, .3); - for (int i = 0; i < 90000; i ++){ - int2 pixelcoord = (int2) (i % *grid_width, i / *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); + // 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); + + // x, y + 1 - //write_imagef(front_grid, pixelcoord, black); + int neighbors = 0; - write_imagei(front_grid, pixelcoord, black); + 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); + } } \ No newline at end of file diff --git a/Conway_OpenCL/fragment_shader.sh b/Conway_OpenCL/fragment_shader.sh index 459d9fb..2448770 100644 --- a/Conway_OpenCL/fragment_shader.sh +++ b/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); } \ No newline at end of file