diff options
| -rw-r--r-- | Release/cl/mandelbrot64.cl | 34 | ||||
| -rw-r--r-- | Release/glade/settings.glade | 263 | ||||
| -rw-r--r-- | src/mandelbrot-zoom.c | 10 | ||||
| -rw-r--r-- | src/mandelbrot-zoom.h | 4 | ||||
| -rw-r--r-- | src/render.c | 22 | ||||
| -rw-r--r-- | src/render.h | 1 | ||||
| -rw-r--r-- | src/render_cpu.c | 63 | ||||
| -rw-r--r-- | src/render_cpu.h | 4 | ||||
| -rw-r--r-- | src/render_opencl.c | 89 | ||||
| -rw-r--r-- | src/render_opencl.h | 9 |
10 files changed, 335 insertions, 164 deletions
diff --git a/Release/cl/mandelbrot64.cl b/Release/cl/mandelbrot64.cl index be38b4a..c89f17c 100644 --- a/Release/cl/mandelbrot64.cl +++ b/Release/cl/mandelbrot64.cl @@ -1,11 +1,15 @@ +#define native_log2 log2 +#define native_cos cos + #ifdef KHR_DP_EXTENSION #pragma OPENCL EXTENSION cl_khr_fp64 : enable #else #pragma OPENCL EXTENSION cl_amd_fp64 : enable #endif -__kernel void calculate(__global uchar4 * mandelbrotImage, const double posx, const double posy, const double stepSizeX, const double stepSizeY, const uint maxIterations, const uint width) +__kernel void calculate(__global uchar4 * mandelbrotImage, const double posx, const double posy, const double stepSizeX, const double stepSizeY, const uint maxIterations, const int width) { +//printf("%f, %f, %f, %f, %u, %d", posx, posy, stepSizeX, stepSizeY, maxIterations, width); int tid = get_global_id(0); int i = tid % (width / 4); @@ -162,13 +166,6 @@ __kernel void calculate(__global uchar4 * mandelbrotImage, const double posx, co color[0].s1 = 0; color[0].s2 = 0; } - if (bench) - { - color[0].s0 = ccount.s0 & 0xff; - color[0].s1 = (ccount.s0 & 0xff00) >> 8; - color[0].s2 = (ccount.s0 & 0xff0000) >> 16; - color[0].s3 = (ccount.s0 & 0xff000000) >> 24; - } mandelbrotImage[4 * tid] = color[0]; c = fc.s1 * 2.0 * 3.1416 / 256.0; color[1].s0 = ((1.0 + native_cos(c)) * 0.5) * 255; @@ -181,13 +178,6 @@ __kernel void calculate(__global uchar4 * mandelbrotImage, const double posx, co color[1].s1 = 0; color[1].s2 = 0; } - if (bench) - { - color[1].s0 = ccount.s1 & 0xff; - color[1].s1 = (ccount.s1 & 0xff00) >> 8; - color[1].s2 = (ccount.s1 & 0xff0000) >> 16; - color[1].s3 = (ccount.s1 & 0xff000000) >> 24; - } mandelbrotImage[4 * tid + 1] = color[1]; c = fc.s2 * 2.0 * 3.1416 / 256.0; color[2].s0 = ((1.0 + native_cos(c)) * 0.5) * 255; @@ -200,13 +190,6 @@ __kernel void calculate(__global uchar4 * mandelbrotImage, const double posx, co color[2].s1 = 0; color[2].s2 = 0; } - if (bench) - { - color[2].s0 = ccount.s2 & 0xff; - color[2].s1 = (ccount.s2 & 0xff00) >> 8; - color[2].s2 = (ccount.s2 & 0xff0000) >> 16; - color[2].s3 = (ccount.s2 & 0xff000000) >> 24; - } mandelbrotImage[4 * tid + 2] = color[2]; c = fc.s3 * 2.0 * 3.1416 / 256.0; color[3].s0 = ((1.0 + native_cos(c)) * 0.5) * 255; @@ -219,12 +202,5 @@ __kernel void calculate(__global uchar4 * mandelbrotImage, const double posx, co color[3].s1 = 0; color[3].s2 = 0; } - if (bench) - { - color[3].s0 = ccount.s3 & 0xff; - color[3].s1 = (ccount.s3 & 0xff00) >> 8; - color[3].s2 = (ccount.s3 & 0xff0000) >> 16; - color[3].s3 = (ccount.s3 & 0xff000000) >> 24; - } mandelbrotImage[4 * tid + 3] = color[3]; }
\ No newline at end of file diff --git a/Release/glade/settings.glade b/Release/glade/settings.glade index ff11a2c..8029d00 100644 --- a/Release/glade/settings.glade +++ b/Release/glade/settings.glade @@ -95,7 +95,6 @@ <property name="can_focus">False</property> <property name="n_rows">6</property> <property name="n_columns">2</property> - <property name="homogeneous">True</property> <child> <object class="GtkLabel" id="iterationsLbl"> <property name="visible">True</property> @@ -220,107 +219,6 @@ </packing> </child> <child> - <object class="GtkLabel" id="zoomLbl"> - <property name="visible">True</property> - <property name="can_focus">False</property> - <property name="label" translatable="yes">Zoom to (temp)</property> - </object> - <packing> - <property name="top_attach">3</property> - <property name="bottom_attach">4</property> - </packing> - </child> - <child> - <object class="GtkHBox" id="hbox5"> - <property name="visible">True</property> - <property name="can_focus">False</property> - <child> - <object class="GtkHBox" id="hbox6"> - <property name="visible">True</property> - <property name="can_focus">False</property> - <child> - <object class="GtkLabel" id="zoomToXLbl"> - <property name="visible">True</property> - <property name="can_focus">False</property> - <property name="label" translatable="yes">X</property> - </object> - <packing> - <property name="expand">True</property> - <property name="fill">True</property> - <property name="position">0</property> - </packing> - </child> - <child> - <object class="GtkEntry" id="zoomToXEntry"> - <property name="visible">True</property> - <property name="can_focus">True</property> - <property name="invisible_char">●</property> - <property name="primary_icon_activatable">False</property> - <property name="secondary_icon_activatable">False</property> - <property name="primary_icon_sensitive">True</property> - <property name="secondary_icon_sensitive">True</property> - </object> - <packing> - <property name="expand">True</property> - <property name="fill">True</property> - <property name="position">1</property> - </packing> - </child> - </object> - <packing> - <property name="expand">True</property> - <property name="fill">True</property> - <property name="position">0</property> - </packing> - </child> - <child> - <object class="GtkHBox" id="hbox7"> - <property name="visible">True</property> - <property name="can_focus">False</property> - <child> - <object class="GtkLabel" id="zoomToYLbl"> - <property name="visible">True</property> - <property name="can_focus">False</property> - <property name="label" translatable="yes">Y</property> - </object> - <packing> - <property name="expand">True</property> - <property name="fill">True</property> - <property name="position">0</property> - </packing> - </child> - <child> - <object class="GtkEntry" id="zoomToYEntry"> - <property name="visible">True</property> - <property name="can_focus">True</property> - <property name="invisible_char">●</property> - <property name="primary_icon_activatable">False</property> - <property name="secondary_icon_activatable">False</property> - <property name="primary_icon_sensitive">True</property> - <property name="secondary_icon_sensitive">True</property> - </object> - <packing> - <property name="expand">True</property> - <property name="fill">True</property> - <property name="position">1</property> - </packing> - </child> - </object> - <packing> - <property name="expand">True</property> - <property name="fill">True</property> - <property name="position">1</property> - </packing> - </child> - </object> - <packing> - <property name="left_attach">1</property> - <property name="right_attach">2</property> - <property name="top_attach">3</property> - <property name="bottom_attach">4</property> - </packing> - </child> - <child> <object class="GtkLabel" id="speedLbl"> <property name="visible">True</property> <property name="can_focus">False</property> @@ -420,6 +318,132 @@ <property name="right_attach">2</property> </packing> </child> + <child> + <object class="GtkLabel" id="zoomLbl"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <property name="label" translatable="yes">Zoom</property> + </object> + <packing> + <property name="top_attach">3</property> + <property name="bottom_attach">4</property> + </packing> + </child> + <child> + <object class="GtkFrame" id="zoomFrame"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <property name="label_xalign">0</property> + <child> + <object class="GtkAlignment" id="alignment7"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <property name="left_padding">12</property> + <child> + <object class="GtkVBox" id="vbox3"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <child> + <object class="GtkCheckButton" id="zoomSpecificCb"> + <property name="label" translatable="yes">Zoom to specific coordinates</property> + <property name="visible">True</property> + <property name="can_focus">True</property> + <property name="receives_default">False</property> + <property name="draw_indicator">True</property> + </object> + <packing> + <property name="expand">True</property> + <property name="fill">True</property> + <property name="position">0</property> + </packing> + </child> + <child> + <object class="GtkHBox" id="hbox5"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <child> + <object class="GtkLabel" id="zoomToXLbl"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <property name="label" translatable="yes">X</property> + </object> + <packing> + <property name="expand">True</property> + <property name="fill">True</property> + <property name="position">0</property> + </packing> + </child> + <child> + <object class="GtkEntry" id="zoomToXEntry"> + <property name="visible">True</property> + <property name="can_focus">True</property> + <property name="invisible_char">●</property> + <property name="primary_icon_activatable">False</property> + <property name="secondary_icon_activatable">False</property> + <property name="primary_icon_sensitive">True</property> + <property name="secondary_icon_sensitive">True</property> + </object> + <packing> + <property name="expand">True</property> + <property name="fill">True</property> + <property name="position">1</property> + </packing> + </child> + <child> + <object class="GtkLabel" id="zoomToYLbl"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <property name="label" translatable="yes">Y</property> + </object> + <packing> + <property name="expand">True</property> + <property name="fill">True</property> + <property name="position">2</property> + </packing> + </child> + <child> + <object class="GtkEntry" id="zoomToYEntry"> + <property name="visible">True</property> + <property name="can_focus">True</property> + <property name="invisible_char">●</property> + <property name="primary_icon_activatable">False</property> + <property name="secondary_icon_activatable">False</property> + <property name="primary_icon_sensitive">True</property> + <property name="secondary_icon_sensitive">True</property> + </object> + <packing> + <property name="expand">True</property> + <property name="fill">True</property> + <property name="position">3</property> + </packing> + </child> + </object> + <packing> + <property name="expand">True</property> + <property name="fill">True</property> + <property name="position">1</property> + </packing> + </child> + </object> + </child> + </object> + </child> + <child type="label"> + <object class="GtkLabel" id="zoomLbl1"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <property name="label" translatable="yes">Zoom precision</property> + <property name="use_markup">True</property> + </object> + </child> + </object> + <packing> + <property name="left_attach">1</property> + <property name="right_attach">2</property> + <property name="top_attach">3</property> + <property name="bottom_attach">4</property> + </packing> + </child> </object> </child> </object> @@ -521,7 +545,7 @@ <object class="GtkLabel" id="fpuLbl"> <property name="visible">True</property> <property name="can_focus">False</property> - <property name="label" translatable="yes">FPU float depth</property> + <property name="label" translatable="yes">FPU float precision</property> </object> <packing> <property name="expand">True</property> @@ -863,20 +887,20 @@ <object class="GtkTable" id="table1"> <property name="visible">True</property> <property name="can_focus">False</property> - <property name="n_rows">2</property> + <property name="n_rows">3</property> <property name="n_columns">2</property> <child> <object class="GtkLabel" id="inLbl"> <property name="visible">True</property> <property name="can_focus">False</property> - <property name="label" translatable="yes">Zoom in</property> + <property name="label" translatable="yes">Start zoom</property> </object> </child> <child> - <object class="GtkLabel" id="outLbl"> + <object class="GtkLabel" id="incIterLbl"> <property name="visible">True</property> <property name="can_focus">False</property> - <property name="label" translatable="yes">Zoom out</property> + <property name="label" translatable="yes">Increase iterations</property> </object> <packing> <property name="top_attach">1</property> @@ -884,6 +908,17 @@ </packing> </child> <child> + <object class="GtkLabel" id="decIterLbl"> + <property name="visible">True</property> + <property name="can_focus">False</property> + <property name="label" translatable="yes">Decrease iterations</property> + </object> + <packing> + <property name="top_attach">2</property> + <property name="bottom_attach">3</property> + </packing> + </child> + <child> <object class="GtkComboBoxText" id="inCombo"> <property name="visible">True</property> <property name="can_focus">False</property> @@ -894,7 +929,7 @@ </packing> </child> <child> - <object class="GtkComboBoxText" id="outCombo"> + <object class="GtkComboBoxText" id="incIterCombo"> <property name="visible">True</property> <property name="can_focus">False</property> </object> @@ -905,6 +940,18 @@ <property name="bottom_attach">2</property> </packing> </child> + <child> + <object class="GtkComboBoxText" id="decIterCombo"> + <property name="visible">True</property> + <property name="can_focus">False</property> + </object> + <packing> + <property name="left_attach">1</property> + <property name="right_attach">2</property> + <property name="top_attach">2</property> + <property name="bottom_attach">3</property> + </packing> + </child> </object> </child> </object> diff --git a/src/mandelbrot-zoom.c b/src/mandelbrot-zoom.c index 72b54ba..f009e10 100644 --- a/src/mandelbrot-zoom.c +++ b/src/mandelbrot-zoom.c @@ -22,6 +22,7 @@ int main(int argc, char **argv) ui_settings.threadsSp = GTK_SPIN_BUTTON(gtk_builder_get_object(builder, "threadsSp")); ui_settings.colorFromBtn = GTK_BUTTON(gtk_builder_get_object(builder, "colorFromBtn")); ui_settings.colorToBtn = GTK_BUTTON(gtk_builder_get_object(builder, "colorToBtn")); + ui_settings.zoomSpecificCb = GTK_CHECK_BUTTON(gtk_builder_get_object(builder, "zoomSpecificCb")); 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")); @@ -40,6 +41,9 @@ int main(int argc, char **argv) ui_settings.startBtn = GTK_BUTTON(gtk_builder_get_object(builder, "startBtn")); ui_settings.exitBtn = GTK_BUTTON(gtk_builder_get_object(builder, "exitBtn")); ui_settings.colorDialog = GTK_COLOR_SELECTION_DIALOG(gtk_builder_get_object(builder, "colorDialog")); + ui_settings.inCombo = GTK_COMBO_BOX_TEXT(gtk_builder_get_object(builder, "inCombo")); + ui_settings.incIterCombo = GTK_COMBO_BOX_TEXT(gtk_builder_get_object(builder, "incIterCombo")); + ui_settings.decIterCombo = GTK_COMBO_BOX_TEXT(gtk_builder_get_object(builder, "decIterCombo")); gtk_spin_button_set_increments(ui_settings.iterationsSp, 1, 2); gtk_spin_button_set_increments(ui_settings.threadsSp, 1, 2); @@ -72,6 +76,8 @@ int main(int argc, char **argv) gtk_combo_box_text_append(ui_settings.fpuCombo, NULL, "64 bit"); gtk_combo_box_text_append(ui_settings.fpuCombo, NULL, "128 bit"); + + g_signal_connect(ui_settings.colorFromBtn, "clicked", G_CALLBACK(on_colorFromBtn_clicked), NULL); g_signal_connect(ui_settings.colorToBtn, "clicked", G_CALLBACK(on_colorToBtn_clicked), NULL); g_signal_connect(ui_settings.startBtn, "clicked", G_CALLBACK(on_startBtn_clicked), NULL); @@ -100,14 +106,18 @@ void on_startBtn_clicked() config.config_cpu.iterations = gtk_spin_button_get_value(ui_settings.iterationsSp); config.config_opencl.iterations = gtk_spin_button_get_value(ui_settings.iterationsSp); + config.zoomSpecific = gtk_toggle_button_get_active(GTK_TOGGLE_BUTTON(ui_settings.zoomSpecificCb)); + config.config_cpu.threads = gtk_spin_button_get_value(ui_settings.threadsSp); config.config_opencl.fma = gtk_toggle_button_get_active(GTK_TOGGLE_BUTTON(ui_settings.fmaCb)); config.video = gtk_toggle_button_get_active(GTK_TOGGLE_BUTTON(ui_settings.exportCb)); + config.path = CHAR_PTR(gtk_entry_get_text(ui_settings.exportTf)); + 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); diff --git a/src/mandelbrot-zoom.h b/src/mandelbrot-zoom.h index 8309885..5aa4dc4 100644 --- a/src/mandelbrot-zoom.h +++ b/src/mandelbrot-zoom.h @@ -21,6 +21,7 @@ typedef struct Ui_settings { GtkSpinButton *threadsSp; GtkButton *colorFromBtn; GtkButton *colorToBtn; + GtkCheckButton *zoomSpecificCb; GtkEntry *zoomToXEntry; GtkEntry *zoomToYEntry; GtkEntry *speedEntry; @@ -39,6 +40,9 @@ typedef struct Ui_settings { GtkButton *startBtn; GtkButton *exitBtn; GtkColorSelectionDialog *colorDialog; + GtkComboBoxText *inCombo; + GtkComboBoxText *incIterCombo; + GtkComboBoxText *decIterCombo; } Ui_settings; Ui_settings ui_settings; diff --git a/src/render.c b/src/render.c index 2ebb7c3..6c01029 100644 --- a/src/render.c +++ b/src/render.c @@ -44,12 +44,30 @@ void init_render(Config *config) case MODE_CPU: init_cpu(&_config->config_cpu); glutDisplayFunc(render_cpu); - glutIdleFunc(idle_cpu); + if (_config->zoomSpecific) + { + glutIdleFunc(idle_cpu); + } + else + { + glutIdleFunc(idle_cpu_dummy); + glutKeyboardFunc(keyboard_cpu); + glutMouseFunc(mouse_cpu); + } break; case MODE_OPENCL: init_opencl(&_config->config_opencl); glutDisplayFunc(render_opencl); - glutIdleFunc(idle_opencl); + if (_config->zoomSpecific) + { + glutIdleFunc(idle_opencl); + } + else + { + glutIdleFunc(idle_opencl_dummy); + glutKeyboardFunc(keyboard_opencl); + glutMouseFunc(mouse_opencl); + } break; } diff --git a/src/render.h b/src/render.h index a50027a..52be446 100644 --- a/src/render.h +++ b/src/render.h @@ -17,6 +17,7 @@ typedef struct config { CpuConfig config_cpu; OpenCLConfig config_opencl; u8 mode; + u8 zoomSpecific; u8 video; u8 filetype; u16 width; diff --git a/src/render_cpu.c b/src/render_cpu.c index 5065d4d..64fc9ec 100644 --- a/src/render_cpu.c +++ b/src/render_cpu.c @@ -8,6 +8,15 @@ #include "render_cpu.h" #define HAVE_STRUCT_TIMESPEC +void idle_cpu_dummy(void) +{ +// glBindTexture(GL_TEXTURE_2D, config_cpu->tex); +// glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, config_cpu->width, +// config_cpu->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, config_cpu->arr); +// glBindTexture(GL_TEXTURE_2D, 0); + glutPostRedisplay(); +} + void init_cpu(CpuConfig *config) { config_cpu = config; @@ -21,10 +30,15 @@ void init_cpu(CpuConfig *config) y_max = y_max_s; calculate(x_min, y_min, x_max, y_max, config_cpu->set_func, config_cpu->arr); + t_old_cpu = glutGet(GLUT_ELAPSED_TIME); } void render_cpu(void) { + glBindTexture(GL_TEXTURE_2D, config_cpu->tex); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, config_cpu->width, + config_cpu->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, config_cpu->arr); + glBindTexture(GL_TEXTURE_2D, 0); glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); glBindTexture(GL_TEXTURE_2D, config_cpu->tex); glEnable(GL_TEXTURE_2D); @@ -66,6 +80,9 @@ void calculate_t(void *args) 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); +// color[2].s0 = ((1.0f + native_cos(c)) * 0.5f) * 255; +// color[2].s1 = ((1.0f + native_cos(2.0f * c + 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255; +// color[2].s2 = ((1.0f + native_cos(c - 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255; _args->arr[COORDS(x, y, config_cpu->width)] = (((1<<24)-1)*iterations)/config_cpu->iterations; } } @@ -73,15 +90,14 @@ void calculate_t(void *args) void idle_cpu(void) { - static int t_old; int t = 0, delta = 0; do { t = glutGet(GLUT_ELAPSED_TIME); - delta = t - t_old; + delta = t - t_old_cpu; } while(delta < 16); // TODO: Hardcoded FPS - t_old = t; + t_old_cpu = t; calculate(x_min, y_min, x_max, y_max, config_cpu->set_func, config_cpu->arr); //glGenTextures(1, &tex); @@ -97,3 +113,44 @@ void idle_cpu(void) y_max = y_max_s - config_cpu->zoom_func(ft, (d64)1.0 - config_cpu->to_y); glutPostRedisplay(); } + +void keyboard_cpu(unsigned char key, int mouseX, int mouseY) +{ + switch(key) + { + case 'i': + config_cpu->iterations++; + break; + case 'd': + config_cpu->iterations--; + break; + } +} + +void mouse_cpu(int button, int state, int x, int y) +{ + if (state == GLUT_DOWN) + { + switch (button) + { + case GLUT_LEFT_BUTTON: + if (config_cpu->speed < 0) + config_cpu->speed = (-1) * config_cpu->speed; + break; + case GLUT_RIGHT_BUTTON: + if (config_cpu->speed > 0) + config_cpu->speed = (-1) * config_cpu->speed; + break; + } + config_cpu->to_x = x_min + + ((d64) x * (x_max - x_min)) / config_cpu->width; + config_cpu->to_y = y_min + + ((d64) y * (y_max - y_min)) / config_cpu->height; + t_old_cpu = glutGet(GLUT_ELAPSED_TIME); + glutIdleFunc(idle_cpu); + } + else if (state == GLUT_UP) + { + glutIdleFunc(idle_cpu_dummy); + } +} diff --git a/src/render_cpu.h b/src/render_cpu.h index 5d32661..7469ca2 100644 --- a/src/render_cpu.h +++ b/src/render_cpu.h @@ -49,6 +49,7 @@ typedef struct t_args { CpuConfig *config_cpu; u32 rendercnt; float ft; +int t_old_cpu; d64 x_min, x_max, y_min, y_max; d64 x_min_s, x_max_s, y_min_s, y_max_s; @@ -56,6 +57,9 @@ d64 x_min_s, x_max_s, y_min_s, y_max_s; void init_cpu(CpuConfig *config); void render_cpu(void); void idle_cpu(void); +void idle_cpu_dummy(void); +void keyboard_cpu(unsigned char key, int mouseX, int mouseY); +void mouse_cpu(int button, int state, int x, int y); 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); diff --git a/src/render_opencl.c b/src/render_opencl.c index 84fc314..055a7f3 100644 --- a/src/render_opencl.c +++ b/src/render_opencl.c @@ -7,16 +7,21 @@ #include "render_opencl.h" +void idle_opencl_dummy(void) +{ + glutPostRedisplay(); +} + void init_opencl(OpenCLConfig *config) { - 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; + x_min_s_cl = -2.0; + x_max_s_cl = 1.0; + y_min_s_cl = -1.0; + y_max_s_cl = 1.0; + x_min_cl = x_min_s_cl; + x_max_cl = x_max_s_cl; + y_min_cl = y_min_s_cl; + y_max_cl = y_max_s_cl; config_opencl = config; output = (cl_uint *) malloc((config_opencl->width) * (config_opencl->height) * sizeof(cl_uchar4)); @@ -126,6 +131,7 @@ void init_opencl(OpenCLConfig *config) { flags = strcat(flags, "-D MUL_ADD=mad "); } + printf("flags: %s\n", flags); program = clCreateProgramWithSource(context, 1, (const char **)&cl_src, (const size_t *)&cl_src_sz, &ret); ret = clBuildProgram(program, num_devices, devices, flags, NULL, NULL); for (cl_uint i = 0; i < num_devices; i++) @@ -147,12 +153,12 @@ void render_opencl(void) cl_double y_max_t; cl_float y_max_t_f; - cl_double x_delta = ((x_max - x_min) / (double) config_opencl->width); - cl_double y_delta = -((y_max - y_min) / (double) config_opencl->height); + cl_double x_delta = ((x_max_cl - x_min_cl) / (double) config_opencl->width); + cl_double y_delta = -((y_max_cl - y_min_cl) / (double) config_opencl->height); cl_float x_delta_f = (float) x_delta; cl_float y_delta_f = (float) y_delta; - cl_float x_min_f = (float) x_min; + cl_float x_min_f = (float) x_min_cl; globalThreads[0] = ((config_opencl->width) * (config_opencl->height)) / num_devices; @@ -171,7 +177,7 @@ void render_opencl(void) localThreads[0] = kernelWorkGroupSize; } - y_max_t = (((y_min + y_max) / 2.0) + (y_max - y_min) / 2.0 - ((double) i * (y_max - y_min)) / (double) num_devices); + y_max_t = (((y_min_cl + y_max_cl) / 2.0) + (y_max_cl - y_min_cl) / 2.0 - ((double) i * (y_max_cl - y_min_cl)) / (double) num_devices); ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &outputBuffer[i]); @@ -189,7 +195,7 @@ void render_opencl(void) ret = clSetKernelArg(kernel, 4, sizeof(cl_float), (void *) &y_delta_f); break; case OPENCL_FPU_64: - ret = clSetKernelArg(kernel, 1, sizeof(cl_double), (void *) &x_min); + ret = clSetKernelArg(kernel, 1, sizeof(cl_double), (void *) &x_min_cl); ret = clSetKernelArg(kernel, 2, sizeof(cl_double), (void *) &y_max_t); ret = clSetKernelArg(kernel, 3, sizeof(cl_double), (void *) &x_delta); ret = clSetKernelArg(kernel, 4, sizeof(cl_double), (void *) &y_delta); @@ -238,6 +244,11 @@ void render_opencl(void) ret = clReleaseEvent(events[num_devices - i - 1]); } + glBindTexture(GL_TEXTURE_2D, config_opencl->tex); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, config_opencl->width, + config_opencl->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, config_opencl->arr); + glBindTexture(GL_TEXTURE_2D, 0); + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); glBindTexture(GL_TEXTURE_2D, config_opencl->tex); glEnable(GL_TEXTURE_2D); @@ -256,14 +267,13 @@ void render_opencl(void) void idle_opencl(void) { - static int t_old; int t = 0, delta = 0; do { t = glutGet(GLUT_ELAPSED_TIME); - delta = t - t_old; + delta = t - t_old_opencl; } while (delta < 16); // TODO: Hardcoded FPS - t_old = t; + t_old_opencl = t; //glGenTextures(1, &tex); glBindTexture(GL_TEXTURE_2D, config_opencl->tex); @@ -273,9 +283,48 @@ void idle_opencl(void) config_opencl->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, config_opencl->arr); glBindTexture(GL_TEXTURE_2D, 0); cl_ft += (config_opencl->speed * (delta / 1000.0)); - x_min = x_min_s + config_opencl->zoom_func(cl_ft, (d64) 2.0 + config_opencl->to_x); - y_min = y_min_s + config_opencl->zoom_func(cl_ft, (d64) 1.0 + config_opencl->to_y); - x_max = x_max_s - config_opencl->zoom_func(cl_ft, (d64) 1.0 - config_opencl->to_x); - y_max = y_max_s - config_opencl->zoom_func(cl_ft, (d64) 1.0 - config_opencl->to_y); + x_min_cl = x_min_s_cl + config_opencl->zoom_func(cl_ft, (cl_double) 2.0 + config_opencl->to_x); + y_min_cl = y_min_s_cl + config_opencl->zoom_func(cl_ft, (cl_double) 1.0 + config_opencl->to_y); + x_max_cl = x_max_s_cl - config_opencl->zoom_func(cl_ft, (cl_double) 1.0 - config_opencl->to_x); + y_max_cl = y_max_s_cl - config_opencl->zoom_func(cl_ft, (cl_double) 1.0 - config_opencl->to_y); glutPostRedisplay(); } + +void keyboard_opencl(unsigned char key, int mouseX, int mouseY) +{ + switch (key) + { + case 'i': + config_opencl->iterations++; + break; + case 'd': + config_opencl->iterations--; + break; + } +} + +void mouse_opencl(int button, int state, int x, int y) +{ + if (state == GLUT_DOWN) + { + switch (button) + { + case GLUT_LEFT_BUTTON: + if (config_opencl->speed < 0) + config_opencl->speed = (-1) * config_opencl->speed; + break; + case GLUT_RIGHT_BUTTON: + if (config_opencl->speed > 0) + config_opencl->speed = (-1) * config_opencl->speed; + break; + } + config_opencl->to_x = x_min_cl + ((d64) x * (x_max_cl - x_min_cl)) / config_opencl->width; + config_opencl->to_y = y_min_cl + ((d64) y * (y_max_cl - y_min_cl)) / config_opencl->height; + t_old_opencl = glutGet(GLUT_ELAPSED_TIME); + glutIdleFunc(idle_opencl); + } + else if(state == GLUT_UP) + { + glutIdleFunc(idle_opencl_dummy); + } +} diff --git a/src/render_opencl.h b/src/render_opencl.h index 720678b..bca77ad 100644 --- a/src/render_opencl.h +++ b/src/render_opencl.h @@ -42,8 +42,8 @@ typedef struct config_opencl { OpenCLConfig *config_opencl; -d64 x_min, x_max, y_min, y_max; -d64 x_min_s, x_max_s, y_min_s, y_max_s; +cl_double x_min_cl, x_max_cl, y_min_cl, y_max_cl; +cl_double x_min_s_cl, x_max_s_cl, y_min_s_cl, y_max_s_cl; float cl_ft; cl_uint *output; @@ -58,10 +58,15 @@ cl_mem outputBuffer[MAX_DEVICES]; cl_int width_cl; +int t_old_opencl; + d64 zoom_func(d64 ft, d64 s); void init_opencl(OpenCLConfig *config); void render_opencl(void); void idle_opencl(void); +void idle_opencl_dummy(void); +void keyboard_opencl(unsigned char key, int mouseX, int mouseY); +void mouse_opencl(int button, int state, int x, int y); #endif /* RENDER_OPENCL_H_ */ |