resetreboot/mandelbrot_study
1
0
Fork 0
Code Issues Pull requests Releases Wiki Activity

Added interactive executable...

Browse source
This commit is contained in:
José Carlos Cuevas 2013-09-01 23:22:01 +02:00
parent 9dfb20c051
commit 79f33446fa
3 changed files with 442 additions and 1 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

8
Makefile
View file

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

363
interactive.c Normal file
View file

@ -0,0 +1,363 @@
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <time.h>
#ifdef __APPLE__
#include <OpenCL/opencl.h>
#else
#include <CL/cl.h>
#endif
#include <SDL.h>
#include <SDL_ttf.h>
#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), &current_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), &center_x, 0, NULL, NULL);
ret = clEnqueueWriteBuffer(command_queue, kernel_center_y, CL_TRUE, 0,
sizeof(float), &center_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;
}

72
mandelbrot_inter_kernel.cl Normal file
View file

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