diff --git a/Makefile b/Makefile index a509c22..794d645 100644 --- a/Makefile +++ b/Makefile @@ -18,7 +18,7 @@ LIBS=-lm -lpthread $(SDLLIBS) INCLUDE=-I/usr/include/SDL -I./ -all: mandelclassic clfract test +all: mandelclassic clfract test clfractinteractive mandelclassic: mandel_classic.o $(CC) $(INCLUDE) mandel_classic.o $(LIBS) -o mandelclassic @@ -32,6 +32,12 @@ clfract: clfract.o clfract.o: main.c $(CC) $(CFLAGS) $(INCLUDE) $(LIBS) $(OPENCLLIBS) main.c -o clfract.o +clfractinteractive: clfractinteractive.o + $(CC) $(INCLUDE) clfractinteractive.o $(LIBS) $(OPENCLLIBS) -o clfractinteractive + +clfractinteractive.o: interactive.c + $(CC) $(CFLAGS) $(INCLUDE) $(LIBS) $(OPENCLLIBS) interactive.c -o clfractinteractive.o + test: test.o $(CC) $(INCLUDE) test.o $(LIBS) -o test diff --git a/interactive.c b/interactive.c new file mode 100644 index 0000000..c85977c --- /dev/null +++ b/interactive.c @@ -0,0 +1,363 @@ +#include +#include +#include + +#include + +#ifdef __APPLE__ +#include +#else +#include +#endif + +#include +#include + +#define MAX_SOURCE_SIZE (0x100000) + +float map_x_mandelbrot(int x, int width, float zoom) +{ + return (((float)x / (float)width) * (3.5 * zoom)) - 2.5; +} + +float map_x_julia(int x, int width, float zoom) +{ + return (((float)x / (float)width) * (3.5 * zoom)) - 1.75; +} + +float map_y(int y, int height, float zoom) +{ + return (((float)y / (float)height) * (2.0 * zoom)) - 1.0; +} + +int main(int argn, char **argv) { + + // Init SDL + if(SDL_Init(SDL_INIT_VIDEO) != 0) + fprintf(stderr, "Could not initialize SDL: %s\n", SDL_GetError()); + + printf("SDL Initialized\n"); + + // Create screen surface + SDL_Surface *screen, *message; + int res_x = 800; + int res_y = 600; + int current_line = 0; + int julia_mode = 0; + + if (argn == 1) + { + julia_mode = 0; + } + else if ((argn == 2) && (strcmp(argv[1], "-julia") == 0)) + { + julia_mode = 1; + printf("Julia mode activated...\n"); + } + + screen = SDL_SetVideoMode(res_x, res_y, 0, SDL_DOUBLEBUF); + if(!screen) + fprintf(stderr,"Could not set video mode: %s\n",SDL_GetError()); + + // Set the title bar + SDL_WM_SetCaption("CLFract", "CLFract"); + + //Initialize SDL_ttf + if( TTF_Init() == -1 ) + { + printf("Error setting up TTF module.\n"); + return 1; + } + + // Load a font + TTF_Font *font; + font = TTF_OpenFont("font.ttf", 24); + if (font == NULL) + { + printf("TTF_OpenFont() Failed: %s", TTF_GetError()); + SDL_Quit(); + return 1; + } + + //The color of the font + SDL_Color textColor = { 255, 255, 255 }; + + // Prepare the resolution and sizes and colors... + const int ITERATIONS = 256; + + // Load the kernel source code into the array source_str + FILE *fp; + char *source_str; + size_t source_size; + + if (julia_mode == 0) + fp = fopen("mandelbrot_inter_kernel.cl", "r"); + else + fp = fopen("julia_kernel.cl", "r"); + + if (!fp) { + fprintf(stderr, "Failed to load kernel.\n"); + exit(1); + } + source_str = (char*)malloc(MAX_SOURCE_SIZE); + source_size = fread( source_str, 1, MAX_SOURCE_SIZE, fp); + fclose( fp ); + + // Get platform and device information + cl_platform_id platform_id = NULL; + cl_device_id device_id = NULL; + cl_uint ret_num_devices; + cl_uint ret_num_platforms; + cl_int ret = clGetPlatformIDs(1, &platform_id, &ret_num_platforms); + ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_GPU, 1, + &device_id, &ret_num_devices); + + // Create an OpenCL context + cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret); + + // Create a command queue + cl_command_queue command_queue = clCreateCommandQueue(context, device_id, 0, &ret); + + // Create memory buffers on the device for returning iterations + // 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(int), NULL, &ret); + cl_mem kernel_zoom_level = clCreateBuffer(context, CL_MEM_READ_ONLY, + sizeof(float), NULL, &ret); + cl_mem kernel_center_x = clCreateBuffer(context, CL_MEM_READ_ONLY, + sizeof(float), NULL, &ret); + cl_mem kernel_center_y = 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(int), NULL, &ret); + + // Create a program from the kernel source + cl_program program = clCreateProgramWithSource(context, 1, + (const char **)&source_str, (const size_t *)&source_size, &ret); + + // 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, "fractal_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); + + int graph_line[res_x]; // (int*)malloc(res_x * sizeof(int)); + float zoom = 1.0; // Our current zoom level + float stop_point; + float center_x = 2.5; + float center_y = 1.75; + + if (julia_mode == 0) + stop_point = 0.00001; + else + stop_point = -2.5; + + SDL_Event ev; + int active, motion; + + active = 1; + motion = 0; + + while(active) + { + for (current_line = 0; current_line < res_y; current_line++) + { + // Set the arguments of the kernel + ret = clEnqueueWriteBuffer(command_queue, kernel_current_line, CL_TRUE, 0, + sizeof(int), ¤t_line, 0, NULL, NULL); + + ret = clEnqueueWriteBuffer(command_queue, kernel_zoom_level, CL_TRUE, 0, + sizeof(float), &zoom, 0, NULL, NULL); + + ret = clEnqueueWriteBuffer(command_queue, kernel_center_x, CL_TRUE, 0, + sizeof(float), ¢er_x, 0, NULL, NULL); + + ret = clEnqueueWriteBuffer(command_queue, kernel_center_y, CL_TRUE, 0, + sizeof(float), ¢er_y, 0, NULL, NULL); + + clFinish(command_queue); + + ret = clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *) &kernel_current_line); + ret = clSetKernelArg(kernel, 3, sizeof(cl_mem), (void *) &kernel_zoom_level); + ret = clSetKernelArg(kernel, 5, sizeof(cl_mem), (void *) &kernel_center_x); + ret = clSetKernelArg(kernel, 6, sizeof(cl_mem), (void *) &kernel_center_y); + + // 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); + + 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; + Uint32 *pixel; + // Lock surface + // SDL_LockSurface(screen); + // rank = screen->pitch/sizeof(Uint32); + pixel = (Uint32*)screen->pixels; + int iteration; + + 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; + // Get the iterations for the point + // printf("Point %d\n", i); + iteration = graph_line[line_count]; + if ((iteration < 128) && (iteration > 0)) { + pixel[(current_line * res_x) + line_count] = SDL_MapRGBA(screen->format, + 0, + 20 + iteration, + 0, + 255); + } + else if ((iteration >= 128) && (iteration < ITERATIONS)) + { + pixel[(current_line * res_x) + line_count] = SDL_MapRGBA(screen->format, + iteration, + 148, + iteration, + 255); + } + else + { + pixel[(current_line * res_x) + line_count] = SDL_MapRGBA(screen->format, + 0, + 0, + 0, + 255); + } + } + } + + // Step, iterate our zoom levels if we're doing mandelbrot or julia set + /* Handle events */ + while(SDL_PollEvent(&ev)) + { + if(ev.type == SDL_QUIT) + active = 0; /* End */ + + else if (ev.type == SDL_MOUSEBUTTONDOWN) + { + SDL_MouseButtonEvent button = ev.button; + if ( (button.state == SDL_PRESSED) && (button.button == SDL_BUTTON_LEFT) ) + { + motion = 1; + } + else if ( (button.state == SDL_PRESSED) && (button.button == SDL_BUTTON_RIGHT) ) + { + motion = -1; + } + + if (julia_mode == 0) + center_x = map_x_mandelbrot(button.x, res_x, zoom); + else + center_x = map_x_julia(button.x, res_x, zoom); + + center_y = map_y(button.y, res_y, zoom); + } + else if ( (ev.type == SDL_MOUSEBUTTONUP) ) + { + motion = 0; + } + } + + if (motion > 0) + { + if (julia_mode == 0) + zoom = zoom * 0.98; + + else + zoom -= 0.01; + } + else if (motion < 0) + { + if (julia_mode == 0) + zoom = zoom / 0.98; + + else + zoom += 0.01; + } + + // Draw message on a corner... + char* msg = (char *)malloc(100 * sizeof(char)); + sprintf(msg, "Zoom level: %0.3f", zoom * 100.0); + message = TTF_RenderText_Solid( font, msg, textColor ); + free(msg); + if (message != NULL) + SDL_BlitSurface(message, NULL, screen, NULL); + + free(message); + + // Draw to the screen + SDL_Flip(screen); + } + + // Clean up + ret = clFlush(command_queue); + ret = clFinish(command_queue); + ret = clReleaseKernel(kernel); + ret = clReleaseProgram(program); + 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_line); + + while(active) + { + } + + SDL_Quit(); + + return 0; +} + diff --git a/mandelbrot_inter_kernel.cl b/mandelbrot_inter_kernel.cl new file mode 100644 index 0000000..86e2209 --- /dev/null +++ b/mandelbrot_inter_kernel.cl @@ -0,0 +1,72 @@ +float map_x(int x, int width, float zoom, float center_x) +{ + return (((float)x / (float)width) * (3.5 * zoom)) - center_x; +} + +float map_y(int y, int height, float zoom, float center_y) +{ + return (((float)y / (float)height) * (2.0 * zoom)) - center_y; +} + +__kernel void fractal_point(__global const int *res_x, + __global const int *res_y, + __global const int *line, + __global const float *zoom, + __global int *graph_line, + __global float *center_x, + __global float *center_y) +{ + // 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, *zoom, *center_x); + float pos_y = map_y(image_y, *res_y, *zoom, *center_y); + float x = 0.0; + float y = 0.0; + float q, x_term; + + // Period-2 bulb check + if (((pos_x + 1.0) * (pos_x + 1.0) + pos_y * pos_y) < 0.0625) + { + graph_line[image_x] = 0; + return; + } + + // Cardioid check + x_term = pos_x - 0.25; + q = x_term * x_term + pos_y * pos_y; + q = q * (q + x_term); + if (q < (0.25 * pos_y * pos_y)) + { + graph_line[image_x] = 0; + return; + } + + int iteration = 0; + int max_iteration = 256; + float xtemp, xx, yy, xplusy; + + while (iteration < max_iteration) + { + xx = x * x; + yy = y * y; + xplusy = x + y; + if ((xx) + (yy) > (4.0)) break; + + xtemp = xx - yy + pos_x; + y = xplusy * xplusy - xx - yy; + y = y + pos_y; + + x = xtemp; + iteration++; + } + + if (iteration > max_iteration) + { + graph_line[image_x] = 0; + } + else + { + graph_line[image_x] = iteration; + } +}