mitchellhansen/conways-game-of-life-gpu
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

Revert

Browse Source
This commit is contained in:
MitchellHansen
2016-01-15 15:07:38 -08:00
parent 7b25b1e845
commit 08fd8b2398
3 changed files with 37 additions and 59 deletions
Download Patch File Download Diff File Expand all files Collapse all files

54
Conway_OpenCL/Conway.hpp
View File

@@ -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, 2); // 25% chance
std::uniform_int_distribution<int> 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);
status = clEnqueueNDRangeKernel(commandQueue, align_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 = 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);

17
Conway_OpenCL/conway_align.cl
View File

@@ -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;
}
}

25
Conway_OpenCL/conway_compute.cl
View File

@@ -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);
}
}