From 1c3ea27064257d8cf7b9f36b0388b8cf2e94ab08 Mon Sep 17 00:00:00 2001 From: Leonard Kugis Date: Fri, 26 Jan 2018 01:12:49 +0100 Subject: Skeleton for OpenCL rendering, cleaned up alot, WIP and crashing --- src/defs.h | 6 +++ src/mandelbrot-zoom.c | 37 ++++++++++++----- src/mandelbrot-zoom.h | 2 + src/mandelbrot.cl | 6 +++ src/render.c | 108 +++++++------------------------------------------- src/render.h | 46 ++++----------------- src/render_cpu.c | 86 ++++++++++++++++++++++++++++++++++++++++ src/render_cpu.h | 58 +++++++++++++++++++++++++++ src/render_opencl.c | 26 ++++++++++++ src/render_opencl.h | 35 ++++++++++++++++ src/sets.c | 14 +++---- src/sets.h | 8 ++-- 12 files changed, 281 insertions(+), 151 deletions(-) create mode 100644 src/mandelbrot.cl create mode 100644 src/render_cpu.c create mode 100644 src/render_cpu.h create mode 100644 src/render_opencl.c create mode 100644 src/render_opencl.h (limited to 'src') diff --git a/src/defs.h b/src/defs.h index 00b51bd..505fe69 100644 --- a/src/defs.h +++ b/src/defs.h @@ -19,6 +19,10 @@ #define FALSE 0 #endif +// Rendering modes, u8 +#define MODE_CPU 0 +#define MODE_OPENCL 1 + typedef unsigned char u8; typedef signed char s8; @@ -31,6 +35,8 @@ typedef signed long int s32; typedef unsigned long long u64; typedef signed long long s64; +typedef long double d64; + //typedef u8 bool; #endif /* DEFS_H_ */ diff --git a/src/mandelbrot-zoom.c b/src/mandelbrot-zoom.c index 9410734..f54a2ab 100644 --- a/src/mandelbrot-zoom.c +++ b/src/mandelbrot-zoom.c @@ -24,6 +24,8 @@ int main(int argc, char **argv) ui_settings.zoomToXEntry = GTK_ENTRY(gtk_builder_get_object(builder, "zoomToXEntry")); ui_settings.zoomToYEntry = GTK_ENTRY(gtk_builder_get_object(builder, "zoomToYEntry")); ui_settings.speedEntry = GTK_ENTRY(gtk_builder_get_object(builder, "speedEntry")); + ui_settings.modeCPURd = GTK_RADIO_BUTTON(gtk_builder_get_object(builder, "modeCPURd")); + ui_settings.modeGPURd = GTK_RADIO_BUTTON(gtk_builder_get_object(builder, "modeGPURd")); ui_settings.exportCb = GTK_CHECK_BUTTON(gtk_builder_get_object(builder, "exportCb")); ui_settings.gifRd = GTK_RADIO_BUTTON(gtk_builder_get_object(builder, "gifRd")); ui_settings.widthSp = GTK_SPIN_BUTTON(gtk_builder_get_object(builder, "widthSp")); @@ -82,12 +84,13 @@ int main(int argc, char **argv) void on_iterationsSp_valueChanged() { - config.iterations = gtk_spin_button_get_value(ui_settings.iterationsSp); + config.config_cpu.iterations = gtk_spin_button_get_value(ui_settings.iterationsSp); + config.config_opencl.iterations = gtk_spin_button_get_value(ui_settings.iterationsSp); } void on_threadsSp_valueChanged() { - config.threads = gtk_spin_button_get_value(ui_settings.threadsSp); + config.config_cpu.threads = gtk_spin_button_get_value(ui_settings.threadsSp); } void on_colorFromBtn_clicked() @@ -115,16 +118,21 @@ void on_exportTf_changed() void on_widthSp_valueChanged() { config.width = gtk_spin_button_get_value(ui_settings.widthSp); + config.config_cpu.width = gtk_spin_button_get_value(ui_settings.widthSp); + config.config_opencl.width = gtk_spin_button_get_value(ui_settings.widthSp); } void on_heightSp_valueChanged() { config.height = gtk_spin_button_get_value(ui_settings.heightSp); + config.config_cpu.height = gtk_spin_button_get_value(ui_settings.heightSp); + config.config_opencl.height = gtk_spin_button_get_value(ui_settings.heightSp); } void on_fpsRenderSp_valueChanged() { - config.renderFPS = gtk_spin_button_get_value(ui_settings.fpsRenderSp); + config.config_cpu.renderFPS = gtk_spin_button_get_value(ui_settings.fpsRenderSp); + config.config_opencl.renderFPS = gtk_spin_button_get_value(ui_settings.fpsRenderSp); } void on_fpsVideoSp_valueChanged() @@ -143,12 +151,23 @@ void on_startBtn_clicked() sscanf(gtk_entry_get_text(ui_settings.zoomToXEntry), "%lf", &x); sscanf(gtk_entry_get_text(ui_settings.zoomToYEntry), "%lf", &y); sscanf(gtk_entry_get_text(ui_settings.speedEntry), "%lf", &speed); - config.to_x = x; - config.to_y = y; - config.speed = speed; - printf("config {\n\t.iterations = %u\n\t.tox = %f\n\t.toy = %f\n\t.video = %u\n\t.filetype = %u\n\t.width = %u\n\t.height = %u\n\t.renderFPS = %u\n\t.videoFPS = %u\n\t.bitrate = %u\n\t.path = %s\n}\n", config.iterations, x, y, config.video, 0, config.width, config.height, config.renderFPS, config.videoFPS, config.bitrate, config.path); - render_init(&config, mandelbrot_r); - render_show(); + config.config_cpu.to_x = x; + config.config_cpu.to_y = y; + config.config_cpu.speed = speed; + config.config_opencl.to_x = x; + config.config_opencl.to_y = y; + config.config_opencl.speed = speed; + //printf("config {\n\t.iterations = %u\n\t.tox = %f\n\t.toy = %f\n\t.video = %u\n\t.filetype = %u\n\t.width = %u\n\t.height = %u\n\t.renderFPS = %u\n\t.videoFPS = %u\n\t.bitrate = %u\n\t.path = %s\n}\n", config.iterations, x, y, config.video, 0, config.width, config.height, config.renderFPS, config.videoFPS, config.bitrate, config.path); + + // TODO: implement selection of sets correctly + config.config_cpu.set_func = mandelbrot_r; + config.config_opencl.set_func = mandelbrot_r; + + // TODO: implement selection of render mode correctly + config.mode = MODE_CPU; + + init_render(&config); + show_render(); } void on_exitBtn_clicked() diff --git a/src/mandelbrot-zoom.h b/src/mandelbrot-zoom.h index a57deba..d30adcf 100644 --- a/src/mandelbrot-zoom.h +++ b/src/mandelbrot-zoom.h @@ -23,6 +23,8 @@ typedef struct Ui_settings { GtkEntry *zoomToXEntry; GtkEntry *zoomToYEntry; GtkEntry *speedEntry; + GtkRadioButton *modeCPURd; + GtkRadioButton *modeGPURd; GtkCheckButton *exportCb; GtkRadioButton *gifRd; GtkSpinButton *widthSp; diff --git a/src/mandelbrot.cl b/src/mandelbrot.cl new file mode 100644 index 0000000..0a477c9 --- /dev/null +++ b/src/mandelbrot.cl @@ -0,0 +1,6 @@ +__kernel void Main(__write_only image2d_t image) +{ + int x = get_global_id(0); + int y = get_global_id(1); + write_imagef(image, (int2)(x, y), (float4)(x / 256.0f, y / 256.0f, 1.0f, 1.0f)); +} \ No newline at end of file diff --git a/src/render.c b/src/render.c index 55d514b..69b46e0 100644 --- a/src/render.c +++ b/src/render.c @@ -9,28 +9,27 @@ #define HAVE_STRUCT_TIMESPEC #include -void render_init(Config *config, u32 (*sfunc) (long double, long double, u32)) +void init_render(Config *config) { - delta = glutGet(GLUT_ELAPSED_TIME); - x_min_s = -2.0; - x_max_s = 1.0; - y_min_s= -1.0; - y_max_s = 1.0; - x_min = x_min_s; - x_max = x_max_s; - y_min = y_min_s; - y_max = y_max_s; - _config = config; - _sfunc = sfunc; s_arr = (u32 *) malloc((_config->width) * (_config->height) * sizeof(u32)); - calculate(x_min, y_min, x_max, y_max, _sfunc, s_arr); //glutInit(0, NULL); glutInitWindowPosition(0, 0); glutInitWindowSize(_config->width, _config->height); glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE | GLUT_DEPTH); glutCreateWindow("Renderer"); - glutDisplayFunc(gl_render); - glutIdleFunc(gl_idle); + + if (_config->mode) + { + glutDisplayFunc(render_cpu); + glutIdleFunc(idle_cpu); + init_cpu(&_config->config_cpu); + } + else + { + glutDisplayFunc(render_opencl); + glutIdleFunc(idle_opencl); + init_opencl(&_config->config_opencl); + } glGenTextures(1, &tex); glBindTexture(GL_TEXTURE_2D, tex); @@ -46,84 +45,7 @@ void render_init(Config *config, u32 (*sfunc) (long double, long double, u32)) glutMainLoop(); } -void render_show() -{ - -} - -void gl_render(void) -{ - glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); - glBindTexture(GL_TEXTURE_2D, tex); - glEnable(GL_TEXTURE_2D); - glBegin(GL_QUADS); - glTexCoord2i(0, 0); glVertex2i(0, 0); - glTexCoord2i(0, 1); glVertex2i(0, _config->height); - glTexCoord2i(1, 1); glVertex2i(_config->width, _config->height); - glTexCoord2i(1, 0); glVertex2i(_config->width, 0); - glEnd(); - glDisable(GL_TEXTURE_2D); - glBindTexture(GL_TEXTURE_2D, 0); - -// glBegin(GL_TRIANGLES); -// glVertex3f(-0.5,-0.5,0.0); -// glVertex3f(0.5,0.0,0.0); -// glVertex3f(0.0,0.5,0.0); -// glEnd(); - - glutSwapBuffers(); -} - -void calculate(long double x_min, long double y_min, long double x_max, long double y_max, u32 (*sfunc) (long double, long double, u32), u32 *arr) +void show_render() { - pthread_t thread; - ThreadArgs *args = (ThreadArgs *) malloc(_config->threads * sizeof(ThreadArgs)); - for(u8 i = 0; i < _config->threads; i++) - { - args[i] = (ThreadArgs) { .tc = _config->threads, .tid = i, .x_min = x_min, .y_min = y_min, .x_max = x_max, .y_max = y_max, .sfunc = sfunc, .arr = arr }; - pthread_create(&thread, NULL, calculate_t, (void *)&args[i]); - } - pthread_join(thread, NULL); -} - -void calculate_t(void *args) -{ - ThreadArgs *_args = (ThreadArgs *)args; - long double x_math, y_math; - u32 iterations; - for (u32 y = (_config->height/_args->tc)*(_args->tid); y < _config->height; y++) - { - for (u32 x = 0; x < _config->width; x++) - { - x_math = _args->x_min + ((long double) x * (_args->x_max - _args->x_min)) / _config->width; - y_math = _args->y_min + ((long double) (_config->height - y) * (_args->y_max - _args->y_min)) / _config->height; - iterations = _args->sfunc(x_math, y_math, _config->iterations); - _args->arr[COORDS(x, y, _config->width)] = (((1<<24)-1)*iterations)/_config->iterations; - } - } -} -long double zoom_func(long double ft, long double s) -{ - return (s - expl(-ft)); -} - -void gl_idle(void) -{ - calculate(x_min, y_min, x_max, y_max, _sfunc, s_arr); - //glGenTextures(1, &tex); - glBindTexture(GL_TEXTURE_2D, tex); - //glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); - //glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); - glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, _config->width, _config->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, s_arr); - glBindTexture(GL_TEXTURE_2D, 0); - int t = glutGet(GLUT_ELAPSED_TIME); - dt = (t - delta) / 1000.0; - delta = t; - ft+=(_config->speed*dt); - x_min = x_min_s + zoom_func(ft, (long double)2.0 + _config->to_x); - y_min = y_min_s + zoom_func(ft, (long double)1.0 + _config->to_y); - x_max = x_max_s - zoom_func(ft, (long double)1.0 - _config->to_x); - y_max = y_max_s - zoom_func(ft, (long double)1.0 - _config->to_y); - glutPostRedisplay(); } diff --git a/src/render.h b/src/render.h index dde79be..2cb76c1 100644 --- a/src/render.h +++ b/src/render.h @@ -11,57 +11,27 @@ #define COORDS(x, y, width) ((y)*(width)+(x)) #include "defs.h" -#include -#include -#include -#include +#include "render_cpu.h" +#include "render_opencl.h" typedef struct config { - u32 iterations; - u8 threads; - u32 colorFrom; - u32 colorTo; - long double to_x; - long double to_y; - long double speed; + CpuConfig config_cpu; + OpenCLConfig config_opencl; + u8 mode; u8 video; u8 filetype; u16 width; u16 height; - u8 renderFPS; u8 videoFPS; u32 bitrate; const char *path; // TODO: key mapping als option in die struct } Config; -typedef struct t_args { - u8 tc; - u8 tid; - long double x_min; - long double x_max; - long double y_min; - long double y_max; - u32 (*sfunc) (long double, long double, u32); - u32 *arr; -} ThreadArgs; - Config *_config; -u32 (*_sfunc) (long double, long double, u32); -u32 *s_arr; -GLuint tex; -u32 rendercnt; -long double x_min, x_max, y_min, y_max; -long double x_min_s, x_max_s, y_min_s, y_max_s; -int delta; -long double dt, ft; - -void render_init(Config *config, u32 (*sfunc) (long double, long double, u32)); -void render_show(); -void gl_render(void); -void gl_idle(void); +u32 s_arr; -void calculate(long double x_min, long double y_min, long double x_max, long double y_max, u32 (*sfunc) (long double, long double, u32), u32 *arr); -void calculate_t(void *args); +void init_render(Config *config); +void show_render(); #endif /* RENDER_H_ */ diff --git a/src/render_cpu.c b/src/render_cpu.c new file mode 100644 index 0000000..9370461 --- /dev/null +++ b/src/render_cpu.c @@ -0,0 +1,86 @@ +/* + * render.c + * + * Created on: 15.01.2018 + * Author: Superleo1810 + */ + +#include "render.h" +#define HAVE_STRUCT_TIMESPEC + +void init_cpu(CpuConfig *config) +{ + config_cpu = config; +} + +void render_cpu(void) +{ + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); + glBindTexture(GL_TEXTURE_2D, tex); + glEnable(GL_TEXTURE_2D); + + glBegin(GL_QUADS); + glTexCoord2i(0, 0); glVertex2i(0, 0); + glTexCoord2i(0, 1); glVertex2i(0, _config->height); + glTexCoord2i(1, 1); glVertex2i(_config->width, _config->height); + glTexCoord2i(1, 0); glVertex2i(_config->width, 0); + glEnd(); + + glDisable(GL_TEXTURE_2D); + glBindTexture(GL_TEXTURE_2D, 0); + glutSwapBuffers(); +} + +void calculate(d64 x_min, d64 y_min, d64 x_max, d64 y_max, u32 (*sfunc) (d64, d64, u32), u32 *arr) +{ + pthread_t thread; + ThreadArgs *args = (ThreadArgs *) malloc(config_cpu->threads * sizeof(ThreadArgs)); + for(u8 i = 0; i < config_cpu->threads; i++) + { + args[i] = (ThreadArgs) { .tc = config_cpu->threads, .tid = i, .x_min = x_min, .y_min = y_min, .x_max = x_max, .y_max = y_max, .sfunc = sfunc, .arr = arr }; + pthread_create(&thread, NULL, calculate_t, (void *)&args[i]); + } + pthread_join(thread, NULL); +} + +void calculate_t(void *args) +{ + ThreadArgs *_args = (ThreadArgs *)args; + d64 x_math, y_math; + u32 iterations; + for (u32 y = (config_cpu->height/_args->tc)*(_args->tid); y < config_cpu->height; y++) + { + for (u32 x = 0; x < _config->width; x++) + { + x_math = _args->x_min + ((d64) x * (_args->x_max - _args->x_min)) / config_cpu->width; + y_math = _args->y_min + ((d64) (config_cpu->height - y) * (_args->y_max - _args->y_min)) / config_cpu->height; + iterations = _args->sfunc(x_math, y_math, config_cpu->iterations); + _args->arr[COORDS(x, y, config_cpu->width)] = (((1<<24)-1)*iterations)/config_cpu->iterations; + } + } +} + +long double zoom_func(d64 ft, d64 s) +{ + return (s - expl(-ft)); +} + +void idle_cpu(void) +{ + calculate(x_min, y_min, x_max, y_max, config_cpu->set_func, config_cpu->arr); + //glGenTextures(1, &tex); + glBindTexture(GL_TEXTURE_2D, tex); + //glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + //glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, config_cpu->width, config_cpu->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, s_arr); + glBindTexture(GL_TEXTURE_2D, 0); + int t = glutGet(GLUT_ELAPSED_TIME); + dt = (t - delta) / 1000.0; + delta = t; + ft+=(config_cpu->speed*dt); + x_min = x_min_s + zoom_func(ft, (d64)2.0 + config_cpu->to_x); + y_min = y_min_s + zoom_func(ft, (d64)1.0 + config_cpu->to_y); + x_max = x_max_s - zoom_func(ft, (d64)1.0 - config_cpu->to_x); + y_max = y_max_s - zoom_func(ft, (d64)1.0 - config_cpu->to_y); + glutPostRedisplay(); +} diff --git a/src/render_cpu.h b/src/render_cpu.h new file mode 100644 index 0000000..6926687 --- /dev/null +++ b/src/render_cpu.h @@ -0,0 +1,58 @@ +/* + * render_cpu.h + * + * Created on: 25.01.2018 + * Author: Superleo1810 + */ + +#ifndef RENDER_CPU_H_ +#define RENDER_CPU_H_ + +#include "defs.h" +#include +#include +#include +#include + +typedef struct config_cpu { + u8 threads; + u32 *arr; + u32 (*set_func)(d64, d64, u32); + u32 iterations; + u32 colorFrom; + u32 colorTo; + d64 to_x; + d64 to_y; + d64 speed; + u16 width; + u16 height; + u8 renderFPS; +} CpuConfig; + +typedef struct t_args { + u8 tc; + u8 tid; + d64 x_min; + d64 x_max; + d64 y_min; + d64 y_max; + u32 (*sfunc) (d64, d64, u32); + u32 *arr; +} ThreadArgs; + + +CpuConfig *config_cpu; +GLuint tex; +u32 rendercnt; +d64 x_min, x_max, y_min, y_max; +d64 x_min_s, x_max_s, y_min_s, y_max_s; +int delta; +d64 dt, ft; + +void init_cpu(CpuConfig *config); +void render_cpu(void); +void idle_cpu(void); +void calculate(d64 x_min, d64 y_min, d64 x_max, d64 y_max, u32 (*sfunc) (d64, d64, u32), u32 *arr); +void calculate_t(void *args); + +#endif /* RENDER_CPU_H_ */ diff --git a/src/render_opencl.c b/src/render_opencl.c new file mode 100644 index 0000000..913c975 --- /dev/null +++ b/src/render_opencl.c @@ -0,0 +1,26 @@ +/* + * render_opencl.c + * + * Created on: 26.01.2018 + * Author: Superleo1810 + */ + +#include "render_opencl.h" + +// TODO: implement opencl + +void init_opencl(OpenCLConfig *config) +{ + config_opencl = config; + printf("OpenCL rendering not implemented yet\n"); +} + +void render_opencl(void) +{ + +} + +void idle_opencl(void) +{ + +} diff --git a/src/render_opencl.h b/src/render_opencl.h new file mode 100644 index 0000000..aaa9df6 --- /dev/null +++ b/src/render_opencl.h @@ -0,0 +1,35 @@ +/* + * render_opencl.h + * + * Created on: 25.01.2018 + * Author: Superleo1810 + */ + +#ifndef RENDER_OPENCL_H_ +#define RENDER_OPENCL_H_ + +#include "defs.h" +#include +#include + +typedef struct config_opencl { + u32 *arr; + u32 (*set_func)(d64, d64, u32); + u32 iterations; + u32 colorFrom; + u32 colorTo; + d64 to_x; + d64 to_y; + d64 speed; + u16 width; + u16 height; + u8 renderFPS; +} OpenCLConfig; + +OpenCLConfig *config_opencl; + +void init_opencl(OpenCLConfig *config); +void render_opencl(void); +void idle_opencl(void); + +#endif /* RENDER_OPENCL_H_ */ diff --git a/src/sets.c b/src/sets.c index 2c01987..86c68f4 100644 --- a/src/sets.c +++ b/src/sets.c @@ -7,9 +7,9 @@ #include "sets.h" -u32 mandelbrot_s(long double x, long double y, u32 iterations) +u32 mandelbrot_s(d64 x, d64 y, u32 iterations) { - long double cx = x, cy = y, x2; + d64 cx = x, cy = y, x2; u32 m = 0; while(m <= iterations && (x*x)+(y*y) <= 4) { @@ -21,16 +21,16 @@ u32 mandelbrot_s(long double x, long double y, u32 iterations) return m; } -u32 mandelbrot_r(long double x, long double y, u32 iterations) +u32 mandelbrot_r(d64 x, d64 y, u32 iterations) { return _mandelbrot_r(x, y, 0.0, 0.0, 0, iterations, 4.0); } -u32 _mandelbrot_r(long double x, long double y, long double zx, long double zy, u32 n, u32 iterations, long double threshold) +u32 _mandelbrot_r(d64 x, d64 y, d64 zx, d64 zy, u32 n, u32 iterations, d64 threshold) { if ((n < iterations) && ((zx * zx + zy * zy) < threshold)) { - long double zx_new = (zx * zx - zy * zy + x); - long double zy_new = (2 * zx * zy + y); + d64 zx_new = (zx * zx - zy * zy + x); + d64 zy_new = (2 * zx * zy + y); if ((zx_new == zx) && (zy_new == zy)) { return iterations; } @@ -39,7 +39,7 @@ u32 _mandelbrot_r(long double x, long double y, long double zx, long double zy, return n; } -u32 julia(long double x, long double y, u32 iterations) +u32 julia(d64 x, d64 y, u32 iterations) { // TODO: Julia-Menge return 0; diff --git a/src/sets.h b/src/sets.h index 1ae237c..cea87fc 100644 --- a/src/sets.h +++ b/src/sets.h @@ -10,9 +10,9 @@ #include "defs.h" -u32 mandelbrot_s(long double x, long double y, u32 iterations); -u32 mandelbrot_r(long double x, long double y, u32 iterations); -u32 _mandelbrot_r(long double x, long double y, long double zx, long double zy, u32 n, u32 iterations, long double threshold); -u32 julia(long double x, long double y, u32 iterations); +u32 mandelbrot_s(d64 x, d64 y, u32 iterations); +u32 mandelbrot_r(d64 x, d64 y, u32 iterations); +u32 _mandelbrot_r(d64 x, d64 y, d64 zx, d64 zy, u32 n, u32 iterations, d64 threshold); +u32 julia(d64 x, d64 y, u32 iterations); #endif /* SETS_H_ */ -- cgit v1.2.1