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

So, it is sharing the textures correctly, it is drawing the OpenGL texture fine. But OpenCL will not touch the thing. CodeXL wont debug it for some reason, and I can't read out from enqueueReadBuffer. I really don't know whats going on here

Browse Source
master
MitchellHansen 9 years ago
parent d999828bee
commit d93a57cbdf
3 changed files with 112 additions and 160 deletions
Show Stats Download Patch File Download Diff File

201
Conway_OpenCL/Conway.hpp
Unescape View File

@ -1,4 +1,5 @@
#include <CL/cl.h>
#include <CL/opencl.h>
#include <string.h>
#include <stdio.h>
#include <stdlib.h>
@ -75,6 +76,68 @@ int main(int argc, char* argv[])
int GRID_HEIGHT = WINDOW_Y;
int WORKER_SIZE = 2000;
// ======================================= Setup OpenGL =======================================================
glfwInit();
glfwWindowHint(GLFW_CONTEXT_VERSION_MAJOR, 3);
glfwWindowHint(GLFW_CONTEXT_VERSION_MINOR, 3);
glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE);
glfwWindowHint(GLFW_RESIZABLE, GL_FALSE);
GLFWwindow* gl_window = glfwCreateWindow(GRID_WIDTH, GRID_HEIGHT, "GPU accelerated life", nullptr, nullptr);
glfwMakeContextCurrent(gl_window);
glfwSetKeyCallback(gl_window, key_callback);
glewExperimental = GL_TRUE;
glewInit();
glViewport(0, 0, GRID_WIDTH, GRID_HEIGHT);
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");
GLfloat vertices[] = {
// Positions // Colors // Texture Coords
1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, // Top Right
1.0f, -1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, // Bottom Right
-1.0f, -1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, // Bottom Left
-1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f // Top Left
};
GLuint indices[] = {
0, 1, 3, // First Triangle
1, 2, 3 // Second Triangle
};
GLuint VBO, VAO, EBO;
glGenVertexArrays(1, &VAO);
glGenBuffers(1, &VBO);
glGenBuffers(1, &EBO);
// Bind the Vertex Array Object first, then bind and set vertex buffer(s) and attribute pointer(s).
glBindVertexArray(VAO);
glBindBuffer(GL_ARRAY_BUFFER, VBO);
glBufferData(GL_ARRAY_BUFFER, sizeof(vertices), vertices, GL_STATIC_DRAW);
glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, EBO);
glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(indices), indices, GL_STATIC_DRAW);
// Position attribute
glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(GLfloat), (GLvoid*)0);
glEnableVertexAttribArray(0);
// Color attribute
glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(GLfloat), (GLvoid*)(3 * sizeof(GLfloat)));
glEnableVertexAttribArray(1);
// TexCoord attribute
glVertexAttribPointer(2, 2, GL_FLOAT, GL_FALSE, 8 * sizeof(GLfloat), (GLvoid*)(6 * sizeof(GLfloat)));
glEnableVertexAttribArray(2);
glBindVertexArray(0); // Unbind VAO
// ============================== OpenCL Setup ==================================================================
// Get the platforms
@ -108,22 +171,27 @@ int main(int argc, char* argv[])
devices = new cl_device_id[numDevices];
status = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, numDevices, devices, NULL);
}
HGLRC hGLRC = wglGetCurrentContext();
HDC hDC = wglGetCurrentDC();
cl_context_properties cps[] ={CL_CONTEXT_PLATFORM, (cl_context_properties)platform, CL_GL_CONTEXT_KHR, (cl_context_properties)hGLRC, CL_WGL_HDC_KHR, (cl_context_properties)hDC, 0 };
cl_context context = clCreateContext(NULL,1, devices,NULL,NULL,NULL);
cl_context context = clCreateContext(cps, 1, devices,NULL,NULL,NULL);
cl_command_queue commandQueue = clCreateCommandQueue(context, devices[0], 0, NULL);
// ============================== Kernel Compilation, Setup ====================================================
// Read the kernel from the file to a string
const char *compute_kernel_filename = "conway_compute.cl";
const char *align_kernel_filename = "conway_align.cl";
const char *compute_kernel_filename = "Z:\\VS_Projects\\Conway_OpenCL\\Conway_OpenCL\\conway_compute.cl";
const char *align_kernel_filename = "Z:\\VS_Projects\\Conway_OpenCL\\Conway_OpenCL\\conway_align.cl";
std::string compute_kernel_string;
std::string align_kernel_string;
convertToString(compute_kernel_filename, compute_kernel_string);
convertToString(compute_kernel_filename, align_kernel_string);
convertToString(align_kernel_filename, align_kernel_string);
// Create a program with the source
const char *compute_source = compute_kernel_string.c_str();
@ -164,70 +232,9 @@ int main(int argc, char* argv[])
}
// Now create the kernels
cl_kernel front_kernel = clCreateKernel(compute_program, "conway_compute", NULL);
cl_kernel compute_kernel = clCreateKernel(compute_program, "conway_compute", NULL);
cl_kernel back_kernel = clCreateKernel(align_program, "conway_align", NULL);
// ======================================= Setup OpenGL =======================================================
glfwInit();
glfwWindowHint(GLFW_CONTEXT_VERSION_MAJOR, 3);
glfwWindowHint(GLFW_CONTEXT_VERSION_MINOR, 3);
glfwWindowHint(GLFW_OPENGL_PROFILE, GLFW_OPENGL_CORE_PROFILE);
glfwWindowHint(GLFW_RESIZABLE, GL_FALSE);
GLFWwindow* gl_window = glfwCreateWindow(GRID_WIDTH, GRID_HEIGHT, "GPU accelerated life", nullptr, nullptr);
glfwMakeContextCurrent(gl_window);
glfwSetKeyCallback(gl_window, key_callback);
glewExperimental = GL_TRUE;
glewInit();
glViewport(0, 0, GRID_WIDTH, GRID_HEIGHT);
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");
GLfloat vertices[] = {
// Positions // Colors // Texture Coords
1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, // Top Right
1.0f, -1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, // Bottom Right
-1.0f, -1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, // Bottom Left
-1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f // Top Left
};
GLuint indices[] = {
0, 1, 3, // First Triangle
1, 2, 3 // Second Triangle
};
GLuint VBO, VAO, EBO;
glGenVertexArrays(1, &VAO);
glGenBuffers(1, &VBO);
glGenBuffers(1, &EBO);
// Bind the Vertex Array Object first, then bind and set vertex buffer(s) and attribute pointer(s).
glBindVertexArray(VAO);
glBindBuffer(GL_ARRAY_BUFFER, VBO);
glBufferData(GL_ARRAY_BUFFER, sizeof(vertices), vertices, GL_STATIC_DRAW);
glBindBuffer(GL_ELEMENT_ARRAY_BUFFER, EBO);
glBufferData(GL_ELEMENT_ARRAY_BUFFER, sizeof(indices), indices, GL_STATIC_DRAW);
// Position attribute
glVertexAttribPointer(0, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(GLfloat), (GLvoid*)0);
glEnableVertexAttribArray(0);
// Color attribute
glVertexAttribPointer(1, 3, GL_FLOAT, GL_FALSE, 8 * sizeof(GLfloat), (GLvoid*)(3 * sizeof(GLfloat)));
glEnableVertexAttribArray(1);
// TexCoord attribute
glVertexAttribPointer(2, 2, GL_FLOAT, GL_FALSE, 8 * sizeof(GLfloat), (GLvoid*)(6 * sizeof(GLfloat)));
glEnableVertexAttribArray(2);
glBindVertexArray(0); // Unbind VAO
// ======================================= Setup grid =========================================================
@ -248,22 +255,16 @@ int main(int argc, char* argv[])
}
}
unsigned char* back_grid = new unsigned char[GRID_WIDTH * GRID_HEIGHT];
for (int i = 0; i < GRID_WIDTH * GRID_HEIGHT; i++) {
back_grid[i] = front_grid[i];
}
// ====================================== Setup Rendering ==========================================================
unsigned char* pixel_array = new sf::Uint8[WINDOW_X * WINDOW_Y * 4];
for (int i = 0; i < GRID_WIDTH * GRID_HEIGHT * 4; i += 4) {
pixel_array[i] = 29; // R?
pixel_array[i] = i % 255; // R?
pixel_array[i + 1] = 70; // G?
pixel_array[i + 2] = 100; // B?
pixel_array[i + 3] = 200; // A?
pixel_array[i + 3] = 100; // A?
}
GLuint texture;
@ -292,45 +293,36 @@ int main(int argc, char* argv[])
glGenerateMipmap(GL_TEXTURE_2D);
delete pixel_array;
//delete pixel_array;
// ========================================= Setup the buffers ==================================================
int err = 0;
cl_mem frontBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, GRID_WIDTH * GRID_HEIGHT * sizeof(char), (void*)front_grid, &err);
cl_mem backBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, GRID_WIDTH * GRID_HEIGHT * sizeof(char), (void*)back_grid, &err);
//cl_mem pixelBuffer = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, GRID_WIDTH * GRID_HEIGHT * sizeof(char), (void*)pixel_array, &err);
cl_mem frontBuffer = clCreateFromGLTexture(context , CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, texture, &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 for front kernel
status = clSetKernelArg(front_kernel, 0, sizeof(cl_mem), (void *)&frontBuffer);
status = clSetKernelArg(front_kernel, 1, sizeof(cl_mem), (void *)&backBuffer);
//status = clSetKernelArg(front_kernel, 2, sizeof(cl_mem), (void *)&pixelBuffer);
status = clSetKernelArg(front_kernel, 3, sizeof(cl_mem), (void *)&workerCountBuffer);
status = clSetKernelArg(front_kernel, 4, sizeof(cl_mem), (void *)&gridWidthBuffer);
status = clSetKernelArg(front_kernel, 5, sizeof(cl_mem), (void *)&gridHeightBuffer);
// Flipped kernel args for the back kernel
status = clSetKernelArg(back_kernel, 0, sizeof(cl_mem), (void *)&backBuffer); // Flipped
status = clSetKernelArg(back_kernel, 1, sizeof(cl_mem), (void *)&frontBuffer); // Flipped
//status = clSetKernelArg(back_kernel, 2, sizeof(cl_mem), (void *)&pixelBuffer);
// 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);
status = clSetKernelArg(back_kernel, 3, sizeof(cl_mem), (void *)&workerCountBuffer);
status = clSetKernelArg(back_kernel, 4, sizeof(cl_mem), (void *)&gridWidthBuffer);
status = clSetKernelArg(back_kernel, 5, 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);
@ -340,11 +332,15 @@ int main(int argc, char* argv[])
//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] = { WORKER_SIZE };
size_t global_work_size[1] = { 10 };
status = clEnqueueAcquireGLObjects(commandQueue, 1, &frontBuffer, 0, 0, 0);
status = clEnqueueNDRangeKernel(commandQueue, compute_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
status = clEnqueueNDRangeKernel(commandQueue, back_kernel, 1, NULL, global_work_size, NULL, 0, NULL, NULL);
//status = clEnqueueReadBuffer(commandQueue, pixelBuffer, CL_TRUE, 0, GRID_WIDTH * GRID_HEIGHT * 4 * sizeof(unsigned char), (void*)pixel_array, 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);
// ======================================= Rendering Shtuff =================================================
@ -352,9 +348,6 @@ int main(int argc, char* argv[])
glfwPollEvents();
// Render
// Clear the colorbuffer
glClearColor(0.2f, 0.3f, 0.3f, 1.0f);
glClear(GL_COLOR_BUFFER_BIT);
glActiveTexture(GL_TEXTURE0);
glBindTexture(GL_TEXTURE_2D, texture);
@ -377,14 +370,12 @@ int main(int argc, char* argv[])
// Release the buffers
status = clReleaseMemObject(frontBuffer);
status = clReleaseMemObject(backBuffer);
//status = clReleaseMemObject(pixelBuffer);
status = clReleaseMemObject(workerCountBuffer);
status = clReleaseMemObject(gridWidthBuffer);
status = clReleaseMemObject(gridHeightBuffer);
// And the program stuff
status = clReleaseKernel(front_kernel);
status = clReleaseKernel(compute_kernel);
status = clReleaseProgram(compute_program);
status = clReleaseProgram(align_program);
status = clReleaseCommandQueue(commandQueue);

70
Conway_OpenCL/conway_compute.cl
Unescape View File

@ -1,61 +1,21 @@
__kernel void conway_compute(__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_compute(__global unsigned char* front_grid, __global int* num_workers, __global int* grid_width, __global int* grid_height)
{
// Caclulate the start and end range that this worker will be calculating
int num = *grid_width * *grid_height * 4;
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 = 0; i < num ; i += 4){
// x, y + 1
int neighbors = 0;
for (int i = start_range; i < end_range; i++){
// add all 8 blocks to neighbors
neighbors = 0;
// Top
neighbors += front_grid[i - *grid_width];
// Top right
neighbors += front_grid[i - *grid_width + 1];
// Right
neighbors += front_grid[i + 1];
// Bottom Right
neighbors += front_grid[i + *grid_width + 1];
// Bottom
neighbors += front_grid[i + *grid_width];
// Bottom Left
neighbors += front_grid[i + *grid_width - 1];
// Left
neighbors += front_grid[i - 1];
// Top left
neighbors += front_grid[i - *grid_width - 1];
if (neighbors == 3 || (neighbors == 2 && front_grid[i])) {
rear_grid[i] = 1;
pixel_out[i * 4] = 255; // R
pixel_out[i * 4 + 1] = 255; // G
pixel_out[i * 4 + 2] = 255; // B
pixel_out[i * 4 + 3] = 255; // A
}
front_grid[i] = 0;
front_grid[i + 1] = 0;
front_grid[i + 2] = 0;
front_grid[i + 3] = 0;
else {
rear_grid[i] = 0;
pixel_out[i * 4] = 49; // R
pixel_out[i * 4 + 1] = 68; // G
pixel_out[i * 4 + 2] = 72; // B
pixel_out[i * 4 + 3] = 255; // A
}
}
front_grid[90000] = 0;
front_grid[90001] = 0;
front_grid[90002] = 0;
front_grid[90003] = 0;
front_grid[90004] = 0;
front_grid[90005] = 0;
front_grid[90006] = 0;
}

1
Conway_OpenCL/fragment_shader.sh
Unescape View File

@ -10,5 +10,6 @@ uniform sampler2D ourTexture1;
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);
}
Write Preview
Loading…
Cancel
Save