diff --git a/.mandel_classic.c.swp b/.mandel_classic.c.swp index 59426f6..4e914e8 100644 Binary files a/.mandel_classic.c.swp and b/.mandel_classic.c.swp differ diff --git a/Makefile b/Makefile index 4245252..263dce0 100644 --- a/Makefile +++ b/Makefile @@ -5,20 +5,21 @@ CC=gcc SDLFLAGS=$(shell sdl-config --cflags) # Comment this line and uncomment the next to get Julia fractals -# CFLAGS=-c -Wall -O2 $(SDLFLAGS) -CFLAGS=-c -Wall -O2 -DJULIA $(SDLFLAGS) +CFLAGS=-c -Wall -O2 $(SDLFLAGS) +# CFLAGS=-c -Wall -O2 -DJULIA $(SDLFLAGS) # CFLAGS=-c -Wall -ggdb $(SDLFLAGS) # Libs! SDLLIBS=$(shell sdl-config --libs) +OPENCLLIBS=-lOpenCL LIBS=-lm -lpthread $(SDLLIBS) # Includes! -INCLUDE=-I/usr/include/SDL +INCLUDE=-I/usr/include/SDL -I./ -all: mandelclassic test +all: mandelclassic clfract test mandelclassic: mandel_classic.o $(CC) $(INCLUDE) mandel_classic.o $(LIBS) -o mandelclassic @@ -26,6 +27,12 @@ mandelclassic: mandel_classic.o mandelclassic.o: mandel_classic.c $(CC) $(CFLAGS) $(INCLUDE) $(LIBS) mandel_classic.c -o mandel_classic.o +clfract: clfract.o + $(CC) $(INCLUDE) clfract.o $(LIBS) $(OPENCLLIBS) -o clfract + +clfract.o: main.c + $(CC) $(CFLAGS) $(INCLUDE) $(LIBS) $(OPENCLLIBS) main.c -o clfract.o + test: test.o $(CC) $(INCLUDE) test.o $(LIBS) -o test diff --git a/main.c b/main.c index a99d0b6..96f079f 100644 --- a/main.c +++ b/main.c @@ -27,20 +27,13 @@ int main(int argn, char **argv) { int current_line = 0; int total_res = res_x * res_y; - screen = SDL_SetVideoMode(res_x, res_y, 0, SDL_HWSURFACE|SDL_DOUBLEBUF); + screen = SDL_SetVideoMode(res_x, res_y, 0, SDL_DOUBLEBUF); if(!screen) fprintf(stderr,"Could not set video mode: %s\n",SDL_GetError()); // Prepare the resolution and sizes and colors... int i; - int temp; const int ITERATIONS = 256; - int *red_scale = (int*)malloc(sizeof(int)*ITERATIONS); - int *blue_scale = (int*)malloc(sizeof(int)*ITERATIONS); - for(i = 0; i < ITERATIONS; i++) { - red_scale[i] = i; - blue_scale[i] = 255 - i; - } // Load the kernel source code into the array source_str FILE *fp; @@ -72,24 +65,18 @@ int main(int argn, char **argv) { cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret); // Create memory buffers on the device for returning iterations - cl_mem kernel_res_x = clCreateBuffer(context, CL_MEM_READ_ONLY, - sizeof(cl_int), NULL, &ret); - cl_mem kernel_res_y = clCreateBuffer(context, CL_MEM_READ_ONLY, - sizeof(cl_int), NULL, &ret); + // Input parameters + cl_mem kernel_res_x = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int), &res_x, &ret); + cl_mem kernel_res_y = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, + sizeof(int), &res_y, &ret); cl_mem kernel_current_line = clCreateBuffer(context, CL_MEM_READ_ONLY, - sizeof(cl_int), NULL, &ret); + sizeof(int), NULL, &ret); + cl_mem kernel_zoom_level = clCreateBuffer(context, CL_MEM_READ_ONLY, + sizeof(float), NULL, &ret); + // Output buffer cl_mem graph_mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, - res_x * sizeof(cl_int), NULL, &ret); - - // Copy resolution x and y for the kernel - ret = clEnqueueWriteBuffer(command_queue, kernel_res_x, CL_TRUE, 0, - sizeof(cl_int), &res_x, 0, NULL, NULL); - - ret = clEnqueueWriteBuffer(command_queue, kernel_res_y, CL_TRUE, 0, - sizeof(cl_int), &res_y, 0, NULL, NULL); - - ret = clEnqueueWriteBuffer(command_queue, kernel_current_line, CL_TRUE, 0, - sizeof(cl_int), ¤t_line, 0, NULL, NULL); + res_x * sizeof(int), NULL, &ret); // Create a program from the kernel source cl_program program = clCreateProgramWithSource(context, 1, @@ -98,94 +85,140 @@ int main(int argn, char **argv) { // Build the program ret = clBuildProgram(program, 1, &device_id, NULL, NULL, NULL); + // Check if it is correct + printf("clBuildProgram\n"); + cl_build_status build_status; + ret = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_STATUS, sizeof(cl_build_status), &build_status, NULL); + + char *build_log; + size_t ret_val_size; + ret = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, 0, NULL, &ret_val_size); + + build_log = (char *) malloc((ret_val_size + 1) * sizeof(char)); + ret = clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG, ret_val_size, build_log, NULL); + build_log[ret_val_size] = '\0'; + printf("BUILD LOG: \n %s", build_log); + printf("program built\n"); + // Create the OpenCL kernel cl_kernel kernel = clCreateKernel(program, "mandelbrot_point", &ret); + // Common kernel params + ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &kernel_res_x); + ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *) &kernel_res_y); + ret = clSetKernelArg(kernel, 4, sizeof(cl_mem), (void *) &graph_mem_obj); + // Our screen in a linear array - int *graph_dots = (int*)malloc(total_res * sizeof(int)); - cl_int *graph_line = (cl_int*)malloc(res_x * sizeof(cl_int)); + int graph_dots[total_res]; // = (int*)malloc(total_res * sizeof(int)); + int graph_line[res_x]; // (int*)malloc(res_x * sizeof(int)); + float zoom; // Our current zoom level - for (current_line = 0; current_line < 600; current_line++) + for(zoom = 1.0; zoom > 0.00001; zoom = zoom * 0.98) { - // Set the arguments of the kernel - ret = clEnqueueWriteBuffer(command_queue, kernel_current_line, CL_TRUE, 0, - sizeof(cl_int), ¤t_line, 0, NULL, NULL); - - ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), &kernel_res_x); - ret = clSetKernelArg(kernel, 1, sizeof(cl_mem), &kernel_res_y); - ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), &kernel_current_line); - ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), graph_mem_obj); - - // Execute the OpenCL kernel on the list - size_t global_item_size = res_x; // Process the entire screen - size_t local_item_size = 64; // Process in groups of 64 - ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, - &global_item_size, &local_item_size, 0, NULL, NULL); - - // Read the memory buffer graph_mem_obj on the device to the local variable graph_dots - ret = clEnqueueReadBuffer(command_queue, graph_mem_obj, CL_TRUE, 0, - res_x * sizeof(cl_int), graph_line, 0, NULL, NULL); - - for (i = 0; i < 800; i++) + for (current_line = 0; current_line < res_y; current_line++) { - graph_dots[(current_line * 800) + i] = graph_line[i]; - } - } + // Set the arguments of the kernel + ret = clEnqueueWriteBuffer(command_queue, kernel_current_line, CL_TRUE, 0, + sizeof(int), ¤t_line, 0, NULL, NULL); - // Display the result to the screen - /* for(i = 0; i < 3078; i++) - printf("Linear: %d -> %d\n", i, graph_dots[i]); */ + ret = clEnqueueWriteBuffer(command_queue, kernel_zoom_level, CL_TRUE, 0, + sizeof(float), &zoom, 0, NULL, NULL); - printf("Rendering...\n"); + clFinish(command_queue); - int iteration; - Uint32 *pixel; - // Lock surface - SDL_LockSurface(screen); - // rank = screen->pitch/sizeof(Uint32); - pixel = (Uint32*)screen->pixels; - /* Draw all dots */ - for(i = 0;i < total_res;i++) - { - // Get the iterations for the point - // printf("Point %d\n", i); - iteration = graph_dots[i]; - if ((iteration < 1000) && (iteration >= 0)) { - pixel[i] = SDL_MapRGBA(screen->format, - red_scale[iteration], - 0, - blue_scale[iteration], - 255); - } - else - { - pixel[i] = SDL_MapRGBA(screen->format, - 0, - 0, - 0, - 255); - } + ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &kernel_current_line); + ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &kernel_zoom_level); - } - // Unlock surface - SDL_UnlockSurface(screen); + // Execute the OpenCL kernel on the list + size_t global_item_size = res_x; // Process the entire line + size_t local_item_size = 32; // Process in groups of 64 + ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, + &global_item_size, &local_item_size, 0, NULL, NULL); - // Draw to the scree - SDL_Flip(screen); + if (ret != CL_SUCCESS) + { + printf("Error while executing kernel\n"); + printf("Error code %d\n", ret); + } + + // Wait for the computation to finish + clFinish(command_queue); + + // Read the memory buffer graph_mem_obj on the device to the local variable graph_dots + ret = clEnqueueReadBuffer(command_queue, graph_mem_obj, CL_TRUE, 0, + res_x * sizeof(int), graph_line, 0, NULL, NULL); + + if (ret != CL_SUCCESS) + printf("Error while reading results buffer\n"); + + clFinish(command_queue); + + int line_count; + for (line_count = 0; line_count < res_x; line_count++) + { + int temp_val = graph_line[line_count]; + graph_dots[(current_line * res_x) + line_count] = temp_val; + } + } + + int iteration; + Uint32 *pixel; + // Lock surface + // SDL_LockSurface(screen); + // rank = screen->pitch/sizeof(Uint32); + pixel = (Uint32*)screen->pixels; + /* Draw all dots */ + for(i = 0;i < total_res;i++) + { + // Get the iterations for the point + // printf("Point %d\n", i); + iteration = graph_dots[i]; + if ((iteration < 128) && (iteration > 0)) { + pixel[i] = SDL_MapRGBA(screen->format, + 0, + 20 + iteration, + 0, + 255); + } + else if ((iteration >= 128) && (iteration < ITERATIONS)) + { + pixel[i] = SDL_MapRGBA(screen->format, + iteration, + 148, + iteration, + 255); + } + else + { + pixel[i] = SDL_MapRGBA(screen->format, + 0, + 0, + 0, + 255); + } + } + // Unlock surface + // SDL_UnlockSurface(screen); + + // Draw to the scree + SDL_Flip(screen); + } // Clean up ret = clFlush(command_queue); ret = clFinish(command_queue); ret = clReleaseKernel(kernel); ret = clReleaseProgram(program); - // ret = clReleaseMemObject(a_mem_obj); - // ret = clReleaseMemObject(b_mem_obj); + ret = clReleaseMemObject(kernel_res_x); + ret = clReleaseMemObject(kernel_res_y); + ret = clReleaseMemObject(kernel_current_line); ret = clReleaseMemObject(graph_mem_obj); ret = clReleaseCommandQueue(command_queue); ret = clReleaseContext(context); // free(A); // free(B); - free(graph_dots); + // free(graph_dots); + // free(graph_line); SDL_Event ev; int active; diff --git a/mandelbrot_kernel.cl b/mandelbrot_kernel.cl index 029599b..d42f85d 100644 --- a/mandelbrot_kernel.cl +++ b/mandelbrot_kernel.cl @@ -1,39 +1,45 @@ -float map_x(int x, int width) +float map_x(int x, int width, float zoom) { - return (((float)x / (float)width) * 3.5) - 2.5; + return (((float)x / (float)width) * (3.5 * zoom)) - (2.5 - (1.0 - zoom)); } -float map_y(int y, int height) +float map_y(int y, int height, float zoom) { - return (((float)y / (float)height) * 2.0) - 1.0; + return (((float)y / (float)height) * (2.0 * zoom)) - (1.00001 - (1.0 - zoom)); } -__kernel void mandelbrot_point(__global const int res_x, __global const int res_y, __global const int line, __global int *graph_line) +__kernel void mandelbrot_point(__global const int *res_x, + __global const int *res_y, + __global const int *line, + __global const float *zoom, + __global int *graph_line) { // Get the index of the current element int image_x = get_global_id(0); - int image_y = line; - float pos_x = map_x(image_x, res_x); - float pos_y = map_y(image_y, res_y); + int image_y = *line; + float pos_x = map_x(image_x, *res_x, *zoom); + float pos_y = map_y(image_y, *res_y, *zoom); float x = 0.0; float y = 0.0; int iteration = 0; - int max_iteration = 255; - float xtemp; + int max_iteration = 256; + float xtemp, xx, yy; while (iteration < max_iteration) { - xtemp = x * x - y * y + pos_x; + xx = x * x; + yy = y * y; + if ((xx) + (yy) > (4.0)) break; + + xtemp = xx - yy + pos_x; y = 2.0 * x * y + pos_y; x = xtemp; iteration++; - - if ((x * x) + (y * y) >= (4.0)) break; } - if (iteration >= max_iteration) + if (iteration > max_iteration) { graph_line[image_x] = 0; } @@ -41,5 +47,4 @@ __kernel void mandelbrot_point(__global const int res_x, __global const int res_ { graph_line[image_x] = iteration; } - } diff --git a/mandelclassic b/mandelclassic index 3cebd92..4c61ca9 100755 Binary files a/mandelclassic and b/mandelclassic differ diff --git a/test b/test index 3c98183..4db960e 100755 Binary files a/test and b/test differ