OpenCL rendering added, not working properly

19 files changed, 1516 insertions, 108 deletions

diff --git a/.cproject b/.cproject
index 78a622f..6b038e1 100644
--- a/.cproject
+++ b/.cproject
@@ -34,6 +34,9 @@
 								<option defaultValue="gnu.c.optimization.level.none" id="gnu.c.compiler.mingw.exe.debug.option.optimization.level.1870710598" name="Optimization Level" superClass="gnu.c.compiler.mingw.exe.debug.option.optimization.level" useByScannerDiscovery="false" valueType="enumerated"/>
 								<option id="gnu.c.compiler.mingw.exe.debug.option.debugging.level.588576008" name="Debug Level" superClass="gnu.c.compiler.mingw.exe.debug.option.debugging.level" useByScannerDiscovery="false" value="gnu.c.debugging.level.max" valueType="enumerated"/>
 								<option id="gnu.c.compiler.option.include.paths.2103695181" name="Include paths (-I)" superClass="gnu.c.compiler.option.include.paths" useByScannerDiscovery="false" valueType="includePath">
+									<listOptionValue builtIn="false" value="&quot;C:\include&quot;"/>
+									<listOptionValue builtIn="false" value="&quot;${AMDAPPSDKROOT}\include&quot;"/>
+									<listOptionValue builtIn="false" value="&quot;${AMDAPPSDKROOT}\include\CL&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\gtk-3.0&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\gdk-pixbuf-2.0&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\glib-2.0&quot;"/>
@@ -42,8 +45,6 @@
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\lib\glib-2.0\include&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\cairo&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include&quot;"/>
-									<listOptionValue builtIn="false" value="&quot;C:\include&quot;"/>
-									<listOptionValue builtIn="false" value="&quot;${AMDAPPSDKROOT}\include&quot;"/>
 								</option>
 								<inputType id="cdt.managedbuild.tool.gnu.c.compiler.input.480003034" superClass="cdt.managedbuild.tool.gnu.c.compiler.input"/>
 							</tool>
@@ -66,6 +67,8 @@
 									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="OpenCL"/>
 									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="gailutil-3"/>
 									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="pthread"/>
+									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="amdocl"/>
+									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="amdocl12cl"/>
 								</option>
 								<option id="gnu.c.link.option.paths.1322466417" name="Library search path (-L)" superClass="gnu.c.link.option.paths" valueType="libPaths">
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\lib&quot;"/>
@@ -160,6 +163,8 @@
 								<option id="gnu.c.compiler.mingw.exe.release.option.debugging.level.992504606" name="Debug Level" superClass="gnu.c.compiler.mingw.exe.release.option.debugging.level" useByScannerDiscovery="false" value="gnu.c.debugging.level.none" valueType="enumerated"/>
 								<option id="gnu.c.compiler.option.include.paths.2087425696" name="Include paths (-I)" superClass="gnu.c.compiler.option.include.paths" useByScannerDiscovery="false" valueType="includePath">
 									<listOptionValue builtIn="false" value="&quot;C:\include&quot;"/>
+									<listOptionValue builtIn="false" value="&quot;${AMDAPPSDKROOT}\include&quot;"/>
+									<listOptionValue builtIn="false" value="&quot;${AMDAPPSDKROOT}\include\CL&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\gtk-3.0&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\gdk-pixbuf-2.0&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\glib-2.0&quot;"/>
@@ -168,7 +173,6 @@
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\lib\glib-2.0\include&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include\cairo&quot;"/>
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\include&quot;"/>
-									<listOptionValue builtIn="false" value="&quot;${AMDAPPSDKROOT}\include&quot;"/>
 								</option>
 								<option id="gnu.c.compiler.option.misc.other.2093587117" name="Other flags" superClass="gnu.c.compiler.option.misc.other" useByScannerDiscovery="false" value="-c -fmessage-length=0 -mms-bitfields" valueType="string"/>
 								<inputType id="cdt.managedbuild.tool.gnu.c.compiler.input.2079810545" superClass="cdt.managedbuild.tool.gnu.c.compiler.input"/>
@@ -193,6 +197,8 @@
 									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="opengl32"/>
 									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="pthread"/>
 									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="OpenCL"/>
+									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="amdocl"/>
+									<listOptionValue builtIn="false" srcPrefixMapping="" srcRootPath="" value="amdocl12cl"/>
 								</option>
 								<option id="gnu.c.link.option.paths.1055912140" name="Library search path (-L)" superClass="gnu.c.link.option.paths" valueType="libPaths">
 									<listOptionValue builtIn="false" value="&quot;${MSYS_HOME}\mingw32\lib&quot;"/>
diff --git a/.settings/org.eclipse.cdt.core.prefs b/.settings/org.eclipse.cdt.core.prefs
index a8f9202..e4c4db6 100644
--- a/.settings/org.eclipse.cdt.core.prefs
+++ b/.settings/org.eclipse.cdt.core.prefs
@@ -3,3 +3,166 @@ environment/project/cdt.managedbuild.config.gnu.mingw.exe.debug.966253669/append
 environment/project/cdt.managedbuild.config.gnu.mingw.exe.debug.966253669/appendContributed=true
 environment/project/cdt.managedbuild.config.gnu.mingw.exe.release.659598002/append=true
 environment/project/cdt.managedbuild.config.gnu.mingw.exe.release.659598002/appendContributed=true
+org.eclipse.cdt.core.formatter.alignment_for_arguments_in_method_invocation=16
+org.eclipse.cdt.core.formatter.alignment_for_assignment=16
+org.eclipse.cdt.core.formatter.alignment_for_base_clause_in_type_declaration=80
+org.eclipse.cdt.core.formatter.alignment_for_binary_expression=16
+org.eclipse.cdt.core.formatter.alignment_for_compact_if=16
+org.eclipse.cdt.core.formatter.alignment_for_conditional_expression=34
+org.eclipse.cdt.core.formatter.alignment_for_conditional_expression_chain=18
+org.eclipse.cdt.core.formatter.alignment_for_constructor_initializer_list=0
+org.eclipse.cdt.core.formatter.alignment_for_declarator_list=16
+org.eclipse.cdt.core.formatter.alignment_for_enumerator_list=48
+org.eclipse.cdt.core.formatter.alignment_for_expression_list=0
+org.eclipse.cdt.core.formatter.alignment_for_expressions_in_array_initializer=16
+org.eclipse.cdt.core.formatter.alignment_for_member_access=0
+org.eclipse.cdt.core.formatter.alignment_for_overloaded_left_shift_chain=16
+org.eclipse.cdt.core.formatter.alignment_for_parameters_in_method_declaration=16
+org.eclipse.cdt.core.formatter.alignment_for_throws_clause_in_method_declaration=16
+org.eclipse.cdt.core.formatter.brace_position_for_array_initializer=next_line
+org.eclipse.cdt.core.formatter.brace_position_for_block=next_line
+org.eclipse.cdt.core.formatter.brace_position_for_block_in_case=next_line
+org.eclipse.cdt.core.formatter.brace_position_for_method_declaration=next_line
+org.eclipse.cdt.core.formatter.brace_position_for_namespace_declaration=next_line
+org.eclipse.cdt.core.formatter.brace_position_for_switch=next_line
+org.eclipse.cdt.core.formatter.brace_position_for_type_declaration=next_line
+org.eclipse.cdt.core.formatter.comment.line_up_line_comment_in_blocks_on_first_column=false
+org.eclipse.cdt.core.formatter.comment.min_distance_between_code_and_line_comment=1
+org.eclipse.cdt.core.formatter.comment.never_indent_line_comments_on_first_column=true
+org.eclipse.cdt.core.formatter.comment.preserve_white_space_between_code_and_line_comments=true
+org.eclipse.cdt.core.formatter.compact_else_if=true
+org.eclipse.cdt.core.formatter.continuation_indentation=2
+org.eclipse.cdt.core.formatter.continuation_indentation_for_array_initializer=2
+org.eclipse.cdt.core.formatter.format_guardian_clause_on_one_line=false
+org.eclipse.cdt.core.formatter.indent_access_specifier_compare_to_type_header=false
+org.eclipse.cdt.core.formatter.indent_access_specifier_extra_spaces=0
+org.eclipse.cdt.core.formatter.indent_body_declarations_compare_to_access_specifier=true
+org.eclipse.cdt.core.formatter.indent_body_declarations_compare_to_namespace_header=false
+org.eclipse.cdt.core.formatter.indent_breaks_compare_to_cases=true
+org.eclipse.cdt.core.formatter.indent_declaration_compare_to_template_header=false
+org.eclipse.cdt.core.formatter.indent_empty_lines=false
+org.eclipse.cdt.core.formatter.indent_statements_compare_to_block=true
+org.eclipse.cdt.core.formatter.indent_statements_compare_to_body=true
+org.eclipse.cdt.core.formatter.indent_switchstatements_compare_to_cases=true
+org.eclipse.cdt.core.formatter.indent_switchstatements_compare_to_switch=false
+org.eclipse.cdt.core.formatter.indentation.size=4
+org.eclipse.cdt.core.formatter.insert_new_line_after_opening_brace_in_array_initializer=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_after_template_declaration=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_at_end_of_file_if_missing=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_before_catch_in_try_statement=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_before_closing_brace_in_array_initializer=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_before_colon_in_constructor_initializer_list=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_before_else_in_if_statement=insert
+org.eclipse.cdt.core.formatter.insert_new_line_before_identifier_in_function_declaration=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_before_while_in_do_statement=do not insert
+org.eclipse.cdt.core.formatter.insert_new_line_in_empty_block=insert
+org.eclipse.cdt.core.formatter.insert_space_after_assignment_operator=insert
+org.eclipse.cdt.core.formatter.insert_space_after_binary_operator=insert
+org.eclipse.cdt.core.formatter.insert_space_after_closing_angle_bracket_in_template_arguments=insert
+org.eclipse.cdt.core.formatter.insert_space_after_closing_angle_bracket_in_template_parameters=insert
+org.eclipse.cdt.core.formatter.insert_space_after_closing_brace_in_block=insert
+org.eclipse.cdt.core.formatter.insert_space_after_closing_paren_in_cast=insert
+org.eclipse.cdt.core.formatter.insert_space_after_colon_in_base_clause=insert
+org.eclipse.cdt.core.formatter.insert_space_after_colon_in_case=insert
+org.eclipse.cdt.core.formatter.insert_space_after_colon_in_conditional=insert
+org.eclipse.cdt.core.formatter.insert_space_after_colon_in_labeled_statement=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_array_initializer=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_base_types=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_declarator_list=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_enum_declarations=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_expression_list=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_method_declaration_parameters=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_method_declaration_throws=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_method_invocation_arguments=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_template_arguments=insert
+org.eclipse.cdt.core.formatter.insert_space_after_comma_in_template_parameters=insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_angle_bracket_in_template_arguments=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_angle_bracket_in_template_parameters=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_brace_in_array_initializer=insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_bracket=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_cast=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_catch=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_exception_specification=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_for=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_if=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_method_declaration=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_method_invocation=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_parenthesized_expression=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_switch=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_opening_paren_in_while=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_postfix_operator=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_prefix_operator=do not insert
+org.eclipse.cdt.core.formatter.insert_space_after_question_in_conditional=insert
+org.eclipse.cdt.core.formatter.insert_space_after_semicolon_in_for=insert
+org.eclipse.cdt.core.formatter.insert_space_after_unary_operator=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_assignment_operator=insert
+org.eclipse.cdt.core.formatter.insert_space_before_binary_operator=insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_angle_bracket_in_template_arguments=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_angle_bracket_in_template_parameters=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_brace_in_array_initializer=insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_bracket=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_cast=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_catch=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_exception_specification=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_for=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_if=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_method_declaration=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_method_invocation=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_parenthesized_expression=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_switch=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_closing_paren_in_while=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_colon_in_base_clause=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_colon_in_case=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_colon_in_conditional=insert
+org.eclipse.cdt.core.formatter.insert_space_before_colon_in_default=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_colon_in_labeled_statement=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_array_initializer=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_base_types=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_declarator_list=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_enum_declarations=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_expression_list=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_method_declaration_parameters=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_method_declaration_throws=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_method_invocation_arguments=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_template_arguments=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_comma_in_template_parameters=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_angle_bracket_in_template_arguments=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_angle_bracket_in_template_parameters=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_array_initializer=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_block=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_method_declaration=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_namespace_declaration=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_switch=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_brace_in_type_declaration=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_bracket=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_catch=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_exception_specification=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_for=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_if=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_method_declaration=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_method_invocation=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_parenthesized_expression=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_switch=insert
+org.eclipse.cdt.core.formatter.insert_space_before_opening_paren_in_while=insert
+org.eclipse.cdt.core.formatter.insert_space_before_postfix_operator=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_prefix_operator=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_question_in_conditional=insert
+org.eclipse.cdt.core.formatter.insert_space_before_semicolon=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_semicolon_in_for=do not insert
+org.eclipse.cdt.core.formatter.insert_space_before_unary_operator=do not insert
+org.eclipse.cdt.core.formatter.insert_space_between_empty_braces_in_array_initializer=do not insert
+org.eclipse.cdt.core.formatter.insert_space_between_empty_brackets=do not insert
+org.eclipse.cdt.core.formatter.insert_space_between_empty_parens_in_exception_specification=do not insert
+org.eclipse.cdt.core.formatter.insert_space_between_empty_parens_in_method_declaration=do not insert
+org.eclipse.cdt.core.formatter.insert_space_between_empty_parens_in_method_invocation=do not insert
+org.eclipse.cdt.core.formatter.join_wrapped_lines=true
+org.eclipse.cdt.core.formatter.keep_else_statement_on_same_line=false
+org.eclipse.cdt.core.formatter.keep_empty_array_initializer_on_one_line=false
+org.eclipse.cdt.core.formatter.keep_imple_if_on_one_line=false
+org.eclipse.cdt.core.formatter.keep_then_statement_on_same_line=false
+org.eclipse.cdt.core.formatter.lineSplit=80
+org.eclipse.cdt.core.formatter.number_of_empty_lines_to_preserve=1
+org.eclipse.cdt.core.formatter.put_empty_statement_on_new_line=true
+org.eclipse.cdt.core.formatter.tabulation.char=tab
+org.eclipse.cdt.core.formatter.tabulation.size=4
+org.eclipse.cdt.core.formatter.use_tabs_only_for_leading_indentations=false
diff --git a/.settings/org.eclipse.cdt.ui.prefs b/.settings/org.eclipse.cdt.ui.prefs
new file mode 100644
index 0000000..48a50c7
--- /dev/null
+++ b/.settings/org.eclipse.cdt.ui.prefs
@@ -0,0 +1,3 @@
+eclipse.preferences.version=1
+formatter_profile=org.eclipse.cdt.ui.default.allman_profile
+formatter_settings_version=1
diff --git a/Release/cl/mandelbrot32.cl b/Release/cl/mandelbrot32.cl
new file mode 100644
index 0000000..af0d0e9
--- /dev/null
+++ b/Release/cl/mandelbrot32.cl
@@ -0,0 +1,199 @@
+__kernel void calculate(__global uchar4 * mandelbrotImage, const float posx, const float posy, const float stepSizeX, const float stepSizeY, const uint maxIterations, const uint width)
+{
+    int tid = get_global_id(0);
+
+    int i = tid % (width / 4);
+    int j = tid / (width / 4);
+    
+    int4 veci = {4 * i, 4 * i + 1, 4 * i + 2, 4 * i + 3};
+    int4 vecj = {j, j, j, j};
+    
+    float4 x0;
+    x0.s0 = (float)(posx + stepSizeX * (float)veci.s0);
+    x0.s1 = (float)(posx + stepSizeX * (float)veci.s1);
+    x0.s2 = (float)(posx + stepSizeX * (float)veci.s2);
+    x0.s3 = (float)(posx + stepSizeX * (float)veci.s3);
+    float4 y0;
+    y0.s0 = (float)(posy + stepSizeY * (float)vecj.s0);
+    y0.s1 = (float)(posy + stepSizeY * (float)vecj.s1);
+    y0.s2 = (float)(posy + stepSizeY * (float)vecj.s2);
+    y0.s3 = (float)(posy + stepSizeY * (float)vecj.s3);
+
+    float4 x = x0;
+    float4 y = y0;
+    
+    uint iter=0;
+    float4 tmp;
+    int4 stay;
+    int4 ccount = 0;
+
+    stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0f;
+    stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0f;
+    stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0f;
+    stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0f;
+    float4 savx = x;
+    float4 savy = y;
+    for(iter=0; (stay.s0 | stay.s1 | stay.s2 | stay.s3) && (iter < maxIterations); iter+= 16)
+    {
+        x = savx;
+        y = savy;
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+        
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0f;
+        stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0f;
+        stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0f;
+        stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0f;
+
+        savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+        savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+        savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+        savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+        savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+        savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+        savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+        savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        ccount += stay*16;
+    }
+    // Handle remainder
+    if (!(stay.s0 & stay.s1 & stay.s2 & stay.s3))
+    {
+        iter = 16;
+        do
+        {
+            x = savx;
+            y = savy;
+            stay.s0 = ((x.s0 * x.s0 + y.s0 * y.s0) <= 4.0f) && 
+						   (ccount.s0 < maxIterations);
+            stay.s1 = ((x.s1 * x.s1 + y.s1 * y.s1) <= 4.0f) && 
+				       (ccount.s1 < maxIterations);
+            stay.s2 = ((x.s2 * x.s2 + y.s2 * y.s2) <= 4.0f) && 
+				       (ccount.s2 < maxIterations);
+            stay.s3 = ((x.s3 * x.s3 + y.s3 * y.s3) <= 4.0f) && 
+				       (ccount.s3 < maxIterations);
+            tmp = x;
+            x = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+            y = MUL_ADD(2.0f * tmp, y, y0);
+            ccount += stay;
+            iter--;
+            savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+            savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+            savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+            savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+            savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+            savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+            savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+            savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        } while ((stay.s0 | stay.s1 | stay.s2 | stay.s3) && iter);
+    }
+    x = savx;
+    y = savy;
+    float4 fc = convert_float4(ccount);
+    fc.s0 = (float)ccount.s0 + 1 - 
+			           native_log2(native_log2(x.s0 * x.s0 + y.s0 * y.s0));
+    fc.s1 = (float)ccount.s1 + 1 - 
+		           native_log2(native_log2(x.s1 * x.s1 + y.s1 * y.s1));
+    fc.s2 = (float)ccount.s2 + 1 - 
+				   native_log2(native_log2(x.s2 * x.s2 + y.s2 * y.s2));
+    fc.s3 = (float)ccount.s3 + 1 - 
+				   native_log2(native_log2(x.s3 * x.s3 + y.s3 * y.s3));
+				   
+    float c = fc.s0 * 2.0f * 3.1416f / 256.0f;
+    uchar4 color[4];
+    color[0].s0 = ((1.0f + native_cos(c)) * 0.5f) * 255;
+    color[0].s1 = ((1.0f + native_cos(2.0f * c + 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[0].s2 = ((1.0f + native_cos(c - 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[0].s3 = 0xff;
+    if (ccount.s0 == maxIterations)
+    {
+        color[0].s0 = 0;
+        color[0].s1 = 0;
+        color[0].s2 = 0;
+    }
+    mandelbrotImage[4 * tid] = color[0];    
+    c = fc.s1 * 2.0f * 3.1416f / 256.0f;
+    color[1].s0 = ((1.0f + native_cos(c)) * 0.5f) * 255;
+    color[1].s1 = ((1.0f + native_cos(2.0f * c + 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[1].s2 = ((1.0f + native_cos(c - 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[1].s3 = 0xff;
+    if (ccount.s1 == maxIterations)
+    {
+        color[1].s0 = 0;
+        color[1].s1 = 0;
+        color[1].s2 = 0;
+    }
+    mandelbrotImage[4 * tid + 1] = color[1];    
+    c = fc.s2 * 2.0f * 3.1416f / 256.0f;
+    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;
+    color[2].s3 = 0xff;
+    if (ccount.s2 == maxIterations)
+    {
+        color[2].s0 = 0;
+        color[2].s1 = 0;
+        color[2].s2 = 0;
+    }
+    mandelbrotImage[4 * tid + 2] = color[2];    
+    c = fc.s3 * 2.0f * 3.1416f / 256.0f;
+    color[3].s0 = ((1.0f + native_cos(c)) * 0.5f) * 255;
+    color[3].s1 = ((1.0f + native_cos(2.0f * c + 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[3].s2 = ((1.0f + native_cos(c - 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[3].s3 = 0xff;
+    if (ccount.s3 == maxIterations)
+    {
+        color[3].s0 = 0;
+        color[3].s1 = 0;
+        color[3].s2 = 0;
+    }
+    mandelbrotImage[4 * tid + 3] = color[3];
+}
+
+#define native_log2 log2
+#define native_cos  cos
\ No newline at end of file
diff --git a/Release/cl/mandelbrot64.cl b/Release/cl/mandelbrot64.cl
new file mode 100644
index 0000000..be38b4a
--- /dev/null
+++ b/Release/cl/mandelbrot64.cl
@@ -0,0 +1,230 @@
+#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)
+{
+    int tid = get_global_id(0);
+
+    int i = tid % (width / 4);
+    int j = tid / (width / 4);
+    
+    int4 veci = {4 * i, 4 * i + 1, 4 * i + 2, 4 * i + 3};
+    int4 vecj = {j, j, j, j};
+    
+    double4 x0;
+    x0.s0 = (double)(posx + stepSizeX * (double)veci.s0);
+    x0.s1 = (double)(posx + stepSizeX * (double)veci.s1);
+    x0.s2 = (double)(posx + stepSizeX * (double)veci.s2);
+    x0.s3 = (double)(posx + stepSizeX * (double)veci.s3);
+    double4 y0;
+    y0.s0 = (double)(posy + stepSizeY * (double)vecj.s0);
+    y0.s1 = (double)(posy + stepSizeY * (double)vecj.s1);
+    y0.s2 = (double)(posy + stepSizeY * (double)vecj.s2);
+    y0.s3 = (double)(posy + stepSizeY * (double)vecj.s3);
+
+    double4 x = x0;
+    double4 y = y0;
+    
+    uint iter=0;
+    double4 tmp;
+    int4 stay;
+    int4 ccount = 0;
+
+    stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0;
+    stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0;
+    stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0;
+    stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0;
+    double4 savx = x;
+    double4 savy = y;
+    for(iter=0; (stay.s0 | stay.s1 | stay.s2 | stay.s3) && (iter < maxIterations); iter+= 16)
+    {
+        x = savx;
+        y = savy;
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); // tmp = x * x + x0 - y * y;
+        y = MUL_ADD(2.0 * x, y, y0); //y = 2.0 * x * y + y0;
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));// x = tmp * tmp + x0 - y * y;
+        y = MUL_ADD(2.0 * tmp, y, y0); //y = 2.0 * tmp * y + y0;
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0); 
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0;
+        stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0;
+        stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0;
+        stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0;
+
+        savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+        savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+        savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+        savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+        savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+        savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+        savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+        savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        ccount += stay*16;
+    }
+    // Handle remainder
+    if (!(stay.s0 & stay.s1 & stay.s2 & stay.s3))
+    {
+        iter = 16;
+        do
+        {
+            x = savx;
+            y = savy;
+            stay.s0 = ((x.s0 * x.s0 + y.s0 * y.s0) <= 4.0) && 
+						   (ccount.s0 < maxIterations);
+            stay.s1 = ((x.s1 * x.s1 + y.s1 * y.s1) <= 4.0) && 
+				       (ccount.s1 < maxIterations);
+            stay.s2 = ((x.s2 * x.s2 + y.s2 * y.s2) <= 4.0) && 
+				       (ccount.s2 < maxIterations);
+            stay.s3 = ((x.s3 * x.s3 + y.s3 * y.s3) <= 4.0) && 
+				       (ccount.s3 < maxIterations);
+            tmp = x;
+            x = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+            y = MUL_ADD(2.0 * tmp, y, y0); //y = 2.0 * tmp * y + y0;
+            ccount += stay;
+            iter--;
+            savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+            savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+            savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+            savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+            savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+            savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+            savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+            savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        } while ((stay.s0 | stay.s1 | stay.s2 | stay.s3) && iter);
+    }
+    x = savx;
+    y = savy;
+    double4 fc = convert_double4(ccount);
+    fc.s0 = (double)ccount.s0 + 1 - 
+			           native_log2(native_log2(x.s0 * x.s0 + y.s0 * y.s0));
+    fc.s1 = (double)ccount.s1 + 1 - 
+		           native_log2(native_log2(x.s1 * x.s1 + y.s1 * y.s1));
+    fc.s2 = (double)ccount.s2 + 1 - 
+				   native_log2(native_log2(x.s2 * x.s2 + y.s2 * y.s2));
+    fc.s3 = (double)ccount.s3 + 1 - 
+				   native_log2(native_log2(x.s3 * x.s3 + y.s3 * y.s3));
+				   
+    double c = fc.s0 * 2.0 * 3.1416 / 256.0;
+    uchar4 color[4];
+    color[0].s0 = ((1.0 + native_cos(c)) * 0.5) * 255;
+    color[0].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[0].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[0].s3 = 0xff;
+    if (ccount.s0 == maxIterations)
+    {
+        color[0].s0 = 0;
+        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;
+    color[1].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[1].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[1].s3 = 0xff;
+    if (ccount.s1 == maxIterations)
+    {
+        color[1].s0 = 0;
+        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;
+    color[2].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[2].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[2].s3 = 0xff;
+    if (ccount.s2 == maxIterations)
+    {
+        color[2].s0 = 0;
+        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;
+    color[3].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[3].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[3].s3 = 0xff;
+    if (ccount.s3 == maxIterations)
+    {
+        color[3].s0 = 0;
+        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 40b0995..ff11a2c 100644
--- a/Release/glade/settings.glade
+++ b/Release/glade/settings.glade
@@ -356,16 +356,6 @@
                       </object>
                     </child>
                     <child>
-                      <object class="GtkComboBox" id="setCombo">
-                        <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>
-                      </packing>
-                    </child>
-                    <child>
                       <object class="GtkLabel" id="modeLbl">
                         <property name="visible">True</property>
                         <property name="can_focus">False</property>
@@ -388,6 +378,7 @@
                             <property name="receives_default">False</property>
                             <property name="active">True</property>
                             <property name="draw_indicator">True</property>
+                            <property name="group">modeGPURd</property>
                           </object>
                           <packing>
                             <property name="expand">True</property>
@@ -403,6 +394,7 @@
                             <property name="receives_default">False</property>
                             <property name="active">True</property>
                             <property name="draw_indicator">True</property>
+                            <property name="group">modeCPURd</property>
                           </object>
                           <packing>
                             <property name="expand">True</property>
@@ -418,6 +410,16 @@
                         <property name="bottom_attach">6</property>
                       </packing>
                     </child>
+                    <child>
+                      <object class="GtkComboBoxText" id="setCombo">
+                        <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>
+                      </packing>
+                    </child>
                   </object>
                 </child>
               </object>
@@ -508,10 +510,57 @@
                     <property name="can_focus">False</property>
                     <property name="left_padding">12</property>
                     <child>
-                      <object class="GtkLabel" id="label1">
+                      <object class="GtkVBox" id="vbox4">
                         <property name="visible">True</property>
                         <property name="can_focus">False</property>
-                        <property name="label" translatable="yes">not implemented yet</property>
+                        <child>
+                          <object class="GtkHBox" id="hbox10">
+                            <property name="visible">True</property>
+                            <property name="can_focus">False</property>
+                            <child>
+                              <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>
+                              </object>
+                              <packing>
+                                <property name="expand">True</property>
+                                <property name="fill">True</property>
+                                <property name="position">0</property>
+                              </packing>
+                            </child>
+                            <child>
+                              <object class="GtkComboBoxText" id="fpuCombo">
+                                <property name="visible">True</property>
+                                <property name="can_focus">False</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="GtkCheckButton" id="fmaCb">
+                            <property name="label" translatable="yes">Enable Fused Multiply/Add (FMA)</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">1</property>
+                          </packing>
+                        </child>
                       </object>
                     </child>
                   </object>
diff --git a/src/cl/mandelbrot32.cl b/src/cl/mandelbrot32.cl
new file mode 100644
index 0000000..cda933f
--- /dev/null
+++ b/src/cl/mandelbrot32.cl
@@ -0,0 +1,196 @@
+__kernel void calculate(__global uchar4 * mandelbrotImage, const float posx, const float posy, const float stepSizeX, const float stepSizeY, const uint maxIterations, const uint width)
+{
+    int tid = get_global_id(0);
+
+    int i = tid % (width / 4);
+    int j = tid / (width / 4);
+    
+    int4 veci = {4 * i, 4 * i + 1, 4 * i + 2, 4 * i + 3};
+    int4 vecj = {j, j, j, j};
+    
+    float4 x0;
+    x0.s0 = (float)(posx + stepSizeX * (float)veci.s0);
+    x0.s1 = (float)(posx + stepSizeX * (float)veci.s1);
+    x0.s2 = (float)(posx + stepSizeX * (float)veci.s2);
+    x0.s3 = (float)(posx + stepSizeX * (float)veci.s3);
+    float4 y0;
+    y0.s0 = (float)(posy + stepSizeY * (float)vecj.s0);
+    y0.s1 = (float)(posy + stepSizeY * (float)vecj.s1);
+    y0.s2 = (float)(posy + stepSizeY * (float)vecj.s2);
+    y0.s3 = (float)(posy + stepSizeY * (float)vecj.s3);
+
+    float4 x = x0;
+    float4 y = y0;
+    
+    uint iter=0;
+    float4 tmp;
+    int4 stay;
+    int4 ccount = 0;
+
+    stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0f;
+    stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0f;
+    stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0f;
+    stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0f;
+    float4 savx = x;
+    float4 savy = y;
+    for(iter=0; (stay.s0 | stay.s1 | stay.s2 | stay.s3) && (iter < maxIterations); iter+= 16)
+    {
+        x = savx;
+        y = savy;
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+        
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+        y = MUL_ADD(2.0f * x, y, y0);
+        x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0f * tmp, y, y0);
+
+        stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0f;
+        stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0f;
+        stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0f;
+        stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0f;
+
+        savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+        savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+        savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+        savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+        savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+        savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+        savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+        savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        ccount += stay*16;
+    }
+    // Handle remainder
+    if (!(stay.s0 & stay.s1 & stay.s2 & stay.s3))
+    {
+        iter = 16;
+        do
+        {
+            x = savx;
+            y = savy;
+            stay.s0 = ((x.s0 * x.s0 + y.s0 * y.s0) <= 4.0f) && 
+						   (ccount.s0 < maxIterations);
+            stay.s1 = ((x.s1 * x.s1 + y.s1 * y.s1) <= 4.0f) && 
+				       (ccount.s1 < maxIterations);
+            stay.s2 = ((x.s2 * x.s2 + y.s2 * y.s2) <= 4.0f) && 
+				       (ccount.s2 < maxIterations);
+            stay.s3 = ((x.s3 * x.s3 + y.s3 * y.s3) <= 4.0f) && 
+				       (ccount.s3 < maxIterations);
+            tmp = x;
+            x = MUL_ADD(-y, y, MUL_ADD(x, x, x0));
+            y = MUL_ADD(2.0f * tmp, y, y0);
+            ccount += stay;
+            iter--;
+            savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+            savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+            savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+            savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+            savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+            savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+            savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+            savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        } while ((stay.s0 | stay.s1 | stay.s2 | stay.s3) && iter);
+    }
+    x = savx;
+    y = savy;
+    float4 fc = convert_float4(ccount);
+    fc.s0 = (float)ccount.s0 + 1 - 
+			           native_log2(native_log2(x.s0 * x.s0 + y.s0 * y.s0));
+    fc.s1 = (float)ccount.s1 + 1 - 
+		           native_log2(native_log2(x.s1 * x.s1 + y.s1 * y.s1));
+    fc.s2 = (float)ccount.s2 + 1 - 
+				   native_log2(native_log2(x.s2 * x.s2 + y.s2 * y.s2));
+    fc.s3 = (float)ccount.s3 + 1 - 
+				   native_log2(native_log2(x.s3 * x.s3 + y.s3 * y.s3));
+				   
+    float c = fc.s0 * 2.0f * 3.1416f / 256.0f;
+    uchar4 color[4];
+    color[0].s0 = ((1.0f + native_cos(c)) * 0.5f) * 255;
+    color[0].s1 = ((1.0f + native_cos(2.0f * c + 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[0].s2 = ((1.0f + native_cos(c - 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[0].s3 = 0xff;
+    if (ccount.s0 == maxIterations)
+    {
+        color[0].s0 = 0;
+        color[0].s1 = 0;
+        color[0].s2 = 0;
+    }
+    mandelbrotImage[4 * tid] = color[0];    
+    c = fc.s1 * 2.0f * 3.1416f / 256.0f;
+    color[1].s0 = ((1.0f + native_cos(c)) * 0.5f) * 255;
+    color[1].s1 = ((1.0f + native_cos(2.0f * c + 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[1].s2 = ((1.0f + native_cos(c - 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[1].s3 = 0xff;
+    if (ccount.s1 == maxIterations)
+    {
+        color[1].s0 = 0;
+        color[1].s1 = 0;
+        color[1].s2 = 0;
+    }
+    mandelbrotImage[4 * tid + 1] = color[1];    
+    c = fc.s2 * 2.0f * 3.1416f / 256.0f;
+    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;
+    color[2].s3 = 0xff;
+    if (ccount.s2 == maxIterations)
+    {
+        color[2].s0 = 0;
+        color[2].s1 = 0;
+        color[2].s2 = 0;
+    }
+    mandelbrotImage[4 * tid + 2] = color[2];    
+    c = fc.s3 * 2.0f * 3.1416f / 256.0f;
+    color[3].s0 = ((1.0f + native_cos(c)) * 0.5f) * 255;
+    color[3].s1 = ((1.0f + native_cos(2.0f * c + 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[3].s2 = ((1.0f + native_cos(c - 2.0f * 3.1416f / 3.0f)) * 0.5f) * 255;
+    color[3].s3 = 0xff;
+    if (ccount.s3 == maxIterations)
+    {
+        color[3].s0 = 0;
+        color[3].s1 = 0;
+        color[3].s2 = 0;
+    }
+    mandelbrotImage[4 * tid + 3] = color[3];
+}
\ No newline at end of file
diff --git a/src/cl/mandelbrot64.cl b/src/cl/mandelbrot64.cl
new file mode 100644
index 0000000..be38b4a
--- /dev/null
+++ b/src/cl/mandelbrot64.cl
@@ -0,0 +1,230 @@
+#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)
+{
+    int tid = get_global_id(0);
+
+    int i = tid % (width / 4);
+    int j = tid / (width / 4);
+    
+    int4 veci = {4 * i, 4 * i + 1, 4 * i + 2, 4 * i + 3};
+    int4 vecj = {j, j, j, j};
+    
+    double4 x0;
+    x0.s0 = (double)(posx + stepSizeX * (double)veci.s0);
+    x0.s1 = (double)(posx + stepSizeX * (double)veci.s1);
+    x0.s2 = (double)(posx + stepSizeX * (double)veci.s2);
+    x0.s3 = (double)(posx + stepSizeX * (double)veci.s3);
+    double4 y0;
+    y0.s0 = (double)(posy + stepSizeY * (double)vecj.s0);
+    y0.s1 = (double)(posy + stepSizeY * (double)vecj.s1);
+    y0.s2 = (double)(posy + stepSizeY * (double)vecj.s2);
+    y0.s3 = (double)(posy + stepSizeY * (double)vecj.s3);
+
+    double4 x = x0;
+    double4 y = y0;
+    
+    uint iter=0;
+    double4 tmp;
+    int4 stay;
+    int4 ccount = 0;
+
+    stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0;
+    stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0;
+    stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0;
+    stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0;
+    double4 savx = x;
+    double4 savy = y;
+    for(iter=0; (stay.s0 | stay.s1 | stay.s2 | stay.s3) && (iter < maxIterations); iter+= 16)
+    {
+        x = savx;
+        y = savy;
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); // tmp = x * x + x0 - y * y;
+        y = MUL_ADD(2.0 * x, y, y0); //y = 2.0 * x * y + y0;
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));// x = tmp * tmp + x0 - y * y;
+        y = MUL_ADD(2.0 * tmp, y, y0); //y = 2.0 * tmp * y + y0;
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0); 
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        // Two iterations
+        tmp = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+        y = MUL_ADD(2.0 * x, y, y0); 
+		x = MUL_ADD(-y, y, MUL_ADD(tmp, tmp, x0));
+        y = MUL_ADD(2.0 * tmp, y, y0);
+
+        stay.s0 = (x.s0 * x.s0 + y.s0 * y.s0) <= 4.0;
+        stay.s1 = (x.s1 * x.s1 + y.s1 * y.s1) <= 4.0;
+        stay.s2 = (x.s2 * x.s2 + y.s2 * y.s2) <= 4.0;
+        stay.s3 = (x.s3 * x.s3 + y.s3 * y.s3) <= 4.0;
+
+        savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+        savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+        savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+        savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+        savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+        savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+        savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+        savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        ccount += stay*16;
+    }
+    // Handle remainder
+    if (!(stay.s0 & stay.s1 & stay.s2 & stay.s3))
+    {
+        iter = 16;
+        do
+        {
+            x = savx;
+            y = savy;
+            stay.s0 = ((x.s0 * x.s0 + y.s0 * y.s0) <= 4.0) && 
+						   (ccount.s0 < maxIterations);
+            stay.s1 = ((x.s1 * x.s1 + y.s1 * y.s1) <= 4.0) && 
+				       (ccount.s1 < maxIterations);
+            stay.s2 = ((x.s2 * x.s2 + y.s2 * y.s2) <= 4.0) && 
+				       (ccount.s2 < maxIterations);
+            stay.s3 = ((x.s3 * x.s3 + y.s3 * y.s3) <= 4.0) && 
+				       (ccount.s3 < maxIterations);
+            tmp = x;
+            x = MUL_ADD(-y, y, MUL_ADD(x, x, x0)); 
+            y = MUL_ADD(2.0 * tmp, y, y0); //y = 2.0 * tmp * y + y0;
+            ccount += stay;
+            iter--;
+            savx.s0 = (stay.s0 ? x.s0 : savx.s0);
+            savx.s1 = (stay.s1 ? x.s1 : savx.s1);
+            savx.s2 = (stay.s2 ? x.s2 : savx.s2);
+            savx.s3 = (stay.s3 ? x.s3 : savx.s3);
+            savy.s0 = (stay.s0 ? y.s0 : savy.s0);
+            savy.s1 = (stay.s1 ? y.s1 : savy.s1);
+            savy.s2 = (stay.s2 ? y.s2 : savy.s2);
+            savy.s3 = (stay.s3 ? y.s3 : savy.s3);
+        } while ((stay.s0 | stay.s1 | stay.s2 | stay.s3) && iter);
+    }
+    x = savx;
+    y = savy;
+    double4 fc = convert_double4(ccount);
+    fc.s0 = (double)ccount.s0 + 1 - 
+			           native_log2(native_log2(x.s0 * x.s0 + y.s0 * y.s0));
+    fc.s1 = (double)ccount.s1 + 1 - 
+		           native_log2(native_log2(x.s1 * x.s1 + y.s1 * y.s1));
+    fc.s2 = (double)ccount.s2 + 1 - 
+				   native_log2(native_log2(x.s2 * x.s2 + y.s2 * y.s2));
+    fc.s3 = (double)ccount.s3 + 1 - 
+				   native_log2(native_log2(x.s3 * x.s3 + y.s3 * y.s3));
+				   
+    double c = fc.s0 * 2.0 * 3.1416 / 256.0;
+    uchar4 color[4];
+    color[0].s0 = ((1.0 + native_cos(c)) * 0.5) * 255;
+    color[0].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[0].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[0].s3 = 0xff;
+    if (ccount.s0 == maxIterations)
+    {
+        color[0].s0 = 0;
+        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;
+    color[1].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[1].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[1].s3 = 0xff;
+    if (ccount.s1 == maxIterations)
+    {
+        color[1].s0 = 0;
+        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;
+    color[2].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[2].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[2].s3 = 0xff;
+    if (ccount.s2 == maxIterations)
+    {
+        color[2].s0 = 0;
+        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;
+    color[3].s1 = ((1.0 + native_cos(2.0 * c + 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[3].s2 = ((1.0 + native_cos(c - 2.0 * 3.1416 / 3.0)) * 0.5) * 255;
+    color[3].s3 = 0xff;
+    if (ccount.s3 == maxIterations)
+    {
+        color[3].s0 = 0;
+        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/src/defs.h b/src/defs.h
index 505fe69..fc0c91d 100644
--- a/src/defs.h
+++ b/src/defs.h
@@ -8,6 +8,8 @@
 #ifndef DEFS_H_
 #define DEFS_H_
 
+#include <stdio.h>
+
 #ifndef NULL
 #define NULL 0
 #endif
diff --git a/src/mandelbrot-zoom.c b/src/mandelbrot-zoom.c
index f54a2ab..72b54ba 100644
--- a/src/mandelbrot-zoom.c
+++ b/src/mandelbrot-zoom.c
@@ -17,6 +17,7 @@ int main(int argc, char **argv)
 
 	ui_settings.settings = GTK_WINDOW(gtk_builder_get_object(builder, "settings"));
 	//gtk_builder_connect_signals(builder, NULL);
+	ui_settings.setCombo = GTK_COMBO_BOX_TEXT(gtk_builder_get_object(builder, "setCombo"));
 	ui_settings.iterationsSp = GTK_SPIN_BUTTON(gtk_builder_get_object(builder, "iterationsSp"));
 	ui_settings.threadsSp = GTK_SPIN_BUTTON(gtk_builder_get_object(builder, "threadsSp"));
 	ui_settings.colorFromBtn = GTK_BUTTON(gtk_builder_get_object(builder, "colorFromBtn"));
@@ -26,6 +27,8 @@ int main(int argc, char **argv)
 	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.fpuCombo = GTK_COMBO_BOX_TEXT(gtk_builder_get_object(builder, "fpuCombo"));
+	ui_settings.fmaCb = GTK_COMBO_BOX_TEXT(gtk_builder_get_object(builder, "fmaCb"));
 	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"));
@@ -62,17 +65,15 @@ int main(int argc, char **argv)
 	gtk_spin_button_set_range(ui_settings.fpsVideoSp, 1, 60);
 	gtk_spin_button_set_range(ui_settings.bitrateSp, 1, 65535);
 
-	g_signal_connect(ui_settings.iterationsSp, "value-changed", G_CALLBACK(on_iterationsSp_valueChanged), NULL);
-	g_signal_connect(ui_settings.threadsSp, "value-changed", G_CALLBACK(on_threadsSp_valueChanged), NULL);
+	gtk_combo_box_text_append(ui_settings.setCombo, NULL, "Mandelbrot");
+	gtk_combo_box_text_append(ui_settings.setCombo, NULL, "Julia");
+
+	gtk_combo_box_text_append(ui_settings.fpuCombo, NULL, "32 bit");
+	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.exportCb, "toggled", G_CALLBACK(on_exportCb_toggled), NULL);
-	g_signal_connect(ui_settings.widthSp, "value-changed", G_CALLBACK(on_widthSp_valueChanged), NULL);
-	g_signal_connect(ui_settings.heightSp, "value-changed", G_CALLBACK(on_heightSp_valueChanged), NULL);
-	g_signal_connect(ui_settings.fpsRenderSp, "value-changed", G_CALLBACK(on_fpsRenderSp_valueChanged), NULL);
-	g_signal_connect(ui_settings.fpsVideoSp, "value-changed", G_CALLBACK(on_fpsVideoSp_valueChanged), NULL);
-	g_signal_connect(ui_settings.bitrateSp, "value-changed", G_CALLBACK(on_bitrateSp_valueChanged), NULL);
-	g_signal_connect(ui_settings.exportTf, "changed", G_CALLBACK(on_exportTf_changed), NULL);
 	g_signal_connect(ui_settings.startBtn, "clicked", G_CALLBACK(on_startBtn_clicked), NULL);
 	g_signal_connect(ui_settings.exitBtn, "clicked", G_CALLBACK(on_exitBtn_clicked), NULL);
 
@@ -82,17 +83,6 @@ int main(int argc, char **argv)
 	return 0;
 }
 
-void on_iterationsSp_valueChanged()
-{
-	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.config_cpu.threads = gtk_spin_button_get_value(ui_settings.threadsSp);
-}
-
 void on_colorFromBtn_clicked()
 {
 	currentColor = COLOR_FROM;
@@ -105,48 +95,34 @@ void on_colorToBtn_clicked()
 	gtk_dialog_run(ui_settings.colorDialog);
 }
 
-void on_exportCb_toggled()
+void on_startBtn_clicked()
 {
-	config.video = gtk_toggle_button_get_active(ui_settings.exportCb);
-}
+	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.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));
 
-void on_exportTf_changed()
-{
 	config.path = CHAR_PTR(gtk_entry_get_text(ui_settings.exportTf));
-}
 
-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.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()
-{
 	config.videoFPS = gtk_spin_button_get_value(ui_settings.fpsVideoSp);
-}
 
-void on_bitrateSp_valueChanged()
-{
 	config.bitrate = gtk_spin_button_get_value(ui_settings.bitrateSp);
-}
 
-void on_startBtn_clicked()
-{
 	double x, y, speed;
 	sscanf(gtk_entry_get_text(ui_settings.zoomToXEntry), "%lf", &x);
 	sscanf(gtk_entry_get_text(ui_settings.zoomToYEntry), "%lf", &y);
@@ -157,14 +133,33 @@ void on_startBtn_clicked()
 	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;
+	config.mode = gtk_toggle_button_get_active(GTK_TOGGLE_BUTTON(ui_settings.modeGPURd));
+	//config.mode = MODE_CPU;
+
+	if (strcmp(gtk_combo_box_text_get_active_text(ui_settings.setCombo), "Mandelbrot") == 0)
+	{
+		config.config_cpu.set_func = mandelbrot_r;
+		config.config_opencl.set_func = SFUNC_MANDELBROT;
+	}
+	else if (strcmp(gtk_combo_box_text_get_active_text(ui_settings.setCombo), "Julia") == 0)
+	{
+		config.config_cpu.set_func = julia;
+		config.config_opencl.set_func = SFUNC_JULIA;
+	}
+
+	if (strcmp(gtk_combo_box_text_get_active_text(ui_settings.fpuCombo), "32 bit") == 0)
+	{
+		config.config_opencl.fpu = OPENCL_FPU_32;
+	}
+	else if (strcmp(gtk_combo_box_text_get_active_text(ui_settings.fpuCombo), "64 bit") == 0)
+	{
+		config.config_opencl.fpu = OPENCL_FPU_64;
+	}
+	else if (strcmp(gtk_combo_box_text_get_active_text(ui_settings.fpuCombo), "128 bit") == 0)
+	{
+		config.config_opencl.fpu = OPENCL_FPU_128;
+	}
 
 	init_render(&config);
 	show_render();
diff --git a/src/mandelbrot-zoom.h b/src/mandelbrot-zoom.h
index d30adcf..8309885 100644
--- a/src/mandelbrot-zoom.h
+++ b/src/mandelbrot-zoom.h
@@ -16,6 +16,7 @@
 
 typedef struct Ui_settings {
 	GtkWindow *settings;
+	GtkComboBoxText *setCombo;
 	GtkSpinButton *iterationsSp;
 	GtkSpinButton *threadsSp;
 	GtkButton *colorFromBtn;
@@ -25,6 +26,8 @@ typedef struct Ui_settings {
 	GtkEntry *speedEntry;
 	GtkRadioButton *modeCPURd;
 	GtkRadioButton *modeGPURd;
+	GtkComboBoxText *fpuCombo;
+	GtkCheckButton *fmaCb;
 	GtkCheckButton *exportCb;
 	GtkRadioButton *gifRd;
 	GtkSpinButton *widthSp;
diff --git a/src/mandelbrot.cl b/src/mandelbrot.cl
deleted file mode 100644
index 0a477c9..0000000
--- a/src/mandelbrot.cl
+++ /dev/null
@@ -1,6 +0,0 @@
-__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 b9ad5e0..191dcd9 100644
--- a/src/render.c
+++ b/src/render.c
@@ -21,6 +21,24 @@ void init_render(Config *config)
 	glutInitDisplayMode(GLUT_RGBA | GLUT_DOUBLE | GLUT_DEPTH);
 	glutCreateWindow("Renderer");
 
+	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);
+	//glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, _config->width, _config->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, s_arr);
+
+	glMatrixMode(GL_PROJECTION);
+	glOrtho(0, _config->width, 0, _config->height, -1, 1);
+	glMatrixMode(GL_MODELVIEW);
+
+	_config->config_cpu.tex = tex;
+	_config->config_opencl.tex = tex;
+	_config->config_cpu.zoom_func = zoom_func;
+	_config->config_opencl.zoom_func = zoom_func;
+
 	switch(_config->mode)
 	{
 	case MODE_CPU:
@@ -35,17 +53,6 @@ void init_render(Config *config)
 		break;
 	}
 
-	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);
-	//glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, _config->width, _config->height, 0, GL_RGBA, GL_UNSIGNED_BYTE, s_arr);
-
-	glMatrixMode(GL_PROJECTION);
-	glOrtho(0, _config->width, 0, _config->height, -1, 1);
-	glMatrixMode(GL_MODELVIEW);
 	glutMainLoop();
 
 }
@@ -54,3 +61,8 @@ void show_render()
 {
 
 }
+
+d64 zoom_func(d64 ft, d64 s)
+{
+	return (s - expl(-ft));
+}
diff --git a/src/render.h b/src/render.h
index 4ae2360..a50027a 100644
--- a/src/render.h
+++ b/src/render.h
@@ -9,8 +9,9 @@
 #define RENDER_H_
 
 #include "defs.h"
-#include "render_cpu.h"
 #include "render_opencl.h"
+#include "render_cpu.h"
+#include <math.h>
 
 typedef struct config {
 	CpuConfig config_cpu;
@@ -28,6 +29,9 @@ typedef struct config {
 
 Config *_config;
 u32 *s_arr;
+GLuint tex;
+
+d64 zoom_func(d64 ft, d64 s);
 
 void init_render(Config *config);
 void show_render();
diff --git a/src/render_cpu.c b/src/render_cpu.c
index f3bf28e..5065d4d 100644
--- a/src/render_cpu.c
+++ b/src/render_cpu.c
@@ -11,10 +11,9 @@
 void init_cpu(CpuConfig *config)
 {
 	config_cpu = config;
-	delta = glutGet(GLUT_ELAPSED_TIME);
 	x_min_s = -2.0;
 	x_max_s = 1.0;
-	y_min_s= -1.0;
+	y_min_s = -1.0;
 	y_max_s = 1.0;
 	x_min = x_min_s;
 	x_max = x_max_s;
@@ -27,7 +26,7 @@ void init_cpu(CpuConfig *config)
 void render_cpu(void)
 {
 	glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
-	glBindTexture(GL_TEXTURE_2D, tex);
+	glBindTexture(GL_TEXTURE_2D, config_cpu->tex);
 	glEnable(GL_TEXTURE_2D);
 
 	glBegin(GL_QUADS);
@@ -52,6 +51,7 @@ void calculate(d64 x_min, d64 y_min, d64 x_max, d64 y_max, u32 (*sfunc) (d64, d6
 		pthread_create(&thread, NULL, calculate_t, (void *)&args[i]);
 	}
 	pthread_join(thread, NULL);
+	free(args);
 }
 
 void calculate_t(void *args)
@@ -71,27 +71,29 @@ void calculate_t(void *args)
 	}
 }
 
-long double zoom_func(d64 ft, d64 s)
-{
-	return (s - expl(-ft));
-}
-
 void idle_cpu(void)
 {
+	static int t_old;
+	int t = 0, delta = 0;
+	do
+	{
+		t = glutGet(GLUT_ELAPSED_TIME);
+		delta = t - t_old;
+	}
+	while(delta < 16); // TODO: Hardcoded FPS
+	t_old = t;
+
 	calculate(x_min, y_min, x_max, y_max, config_cpu->set_func, config_cpu->arr);
 	//glGenTextures(1, &tex);
-	glBindTexture(GL_TEXTURE_2D, tex);
+	glBindTexture(GL_TEXTURE_2D, config_cpu->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, config_cpu->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);
+	ft+=(config_cpu->speed*(delta/1000.0));
+	x_min = x_min_s + config_cpu->zoom_func(ft, (d64)2.0 + config_cpu->to_x);
+	y_min = y_min_s + config_cpu->zoom_func(ft, (d64)1.0 + config_cpu->to_y);
+	x_max = x_max_s - config_cpu->zoom_func(ft, (d64)1.0 - config_cpu->to_x);
+	y_max = y_max_s - config_cpu->zoom_func(ft, (d64)1.0 - config_cpu->to_y);
 	glutPostRedisplay();
 }
diff --git a/src/render_cpu.h b/src/render_cpu.h
index b4141ae..7abbdc4 100644
--- a/src/render_cpu.h
+++ b/src/render_cpu.h
@@ -12,13 +12,16 @@
 #include <stdlib.h>
 #include <GL/glut.h>
 #include <pthread.h>
+#include <time.h>
 #include <math.h>
 
 #define COORDS(x, y, width) ((y)*(width)+(x))
 
 typedef struct config_cpu {
 	u8 threads;
+	GLuint tex;
 	u32 *arr;
+	d64 (*zoom_func)(d64, d64);
 	u32 (*set_func)(d64, d64, u32);
 	u32 iterations;
 	u32 colorFrom;
@@ -42,14 +45,12 @@ typedef struct t_args {
 	u32 *arr;
 } ThreadArgs;
 
-
 CpuConfig *config_cpu;
-GLuint tex;
 u32 rendercnt;
+float ft;
+
 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);
diff --git a/src/render_opencl.c b/src/render_opencl.c
index 913c975..9d650c2 100644
--- a/src/render_opencl.c
+++ b/src/render_opencl.c
@@ -7,20 +7,291 @@
 
 #include "render_opencl.h"
 
-// TODO: implement opencl
-
 void init_opencl(OpenCLConfig *config)
 {
+	printf("cl init\n");
+	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_opencl = config;
-	printf("OpenCL rendering not implemented yet\n");
+	output = (cl_uint *) malloc((config_opencl->width) * (config_opencl->height) * sizeof(cl_uchar4));
+	context = NULL;
+
+	cl_platform_id platform_id;
+	cl_uint ret_num_devices;
+	cl_uint ret_num_platforms;
+	size_t device_list_size;
+	cl_program program;
+	size_t cl_src_sz;
+
+	clGetPlatformIDs(1, &platform_id, &ret_num_platforms);
+	clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_DEFAULT, 1, &device_id, &ret_num_devices);
+	context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
+	ret = clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &device_list_size);
+	devices = (cl_device_id *)malloc(device_list_size);
+	num_devices = (cl_uint)(device_list_size/sizeof(cl_device_id));
+	clGetContextInfo(context, CL_CONTEXT_DEVICES, device_list_size, devices, NULL);
+
+	printf("lists done\n");
+
+	FILE *fp;
+	char *cl_src, *path, *flags = (char *)malloc(200 * sizeof(char));
+	flags[0] = '\0';
+	switch(config_opencl->fpu)
+	{
+	case OPENCL_FPU_32:
+		switch(config_opencl->set_func)
+		{
+		case SFUNC_JULIA:
+			path = "cl/julia32.cl";
+			break;
+		case SFUNC_MANDELBROT:
+		default:
+			path = "cl/mandelbrot32.cl";
+			break;
+		}
+		break;
+	case OPENCL_FPU_64:
+		switch(config_opencl->set_func)
+		{
+		case SFUNC_JULIA:
+			path = "cl/julia64.cl";
+			break;
+		case SFUNC_MANDELBROT:
+		default:
+			path = "cl/mandelbrot64.cl";
+			break;
+		}
+		int khrFP64 = 0;
+		int amdFP64 = 0;
+		for (cl_uint i = 0; i < num_devices; i++)
+		{
+			char deviceExtensions[8192];
+			ret = clGetDeviceInfo(devices[i], CL_DEVICE_EXTENSIONS,
+					sizeof(deviceExtensions), deviceExtensions, 0);
+			if (strstr(deviceExtensions, "cl_khr_fp64"))
+			{
+				khrFP64++;
+			}
+			else
+			{
+				if (strstr(deviceExtensions, "cl_amd_fp64"))
+				{
+					amdFP64++;
+				}
+			}
+		}
+		if (khrFP64 == num_devices)
+		{
+			flags = strcat(flags, "-D KHR_DP_EXTENSION ");
+		}
+		else if (amdFP64 == num_devices)
+		{
+			flags = strcat(flags, "");
+		}
+		break;
+	case OPENCL_FPU_128:
+		printf("128 bit precision not implemented yet\n");
+		break;
+	}
+
+	cl_src = (char *)malloc(MAX_SOURCE_SIZE * sizeof(char));
+	fp = fopen(path, "r");
+	cl_src_sz = fread(cl_src, 1, MAX_SOURCE_SIZE, fp);
+	fclose(fp);
+
+	printf("reading done\n");
+
+	for (cl_uint i = 0; i < num_devices; i++)
+	{
+		cl_command_queue_properties prop = 0;
+		//			if (sampleArgs->timing)
+		//			{
+		//				prop |= CL_QUEUE_PROFILING_ENABLE;
+		//			}
+		commandQueue[i] = clCreateCommandQueue(context, devices[i], prop, &ret);
+		outputBuffer[i] = clCreateBuffer(context,
+				CL_MEM_WRITE_ONLY | CL_MEM_ALLOC_HOST_PTR,
+				(sizeof(cl_uint) * (config_opencl->width)
+						* (config_opencl->height)) / num_devices, NULL, &ret);
+	}
+	printf("comand puffer\n");
+	if (config_opencl->fma)
+	{
+		flags = strcat(flags, "-D MUL_ADD=fma ");
+	}
+	else
+	{
+		flags = strcat(flags, "-D MUL_ADD=mad ");
+	}
+	printf("flags concat\n");
+	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++)
+	{
+		kernel_vector[i] = clCreateKernel(program, "calculate", &ret);
+	}
+
+	printf("cl init done\n");
+
 }
 
 void render_opencl(void)
 {
+	cl_event events[MAX_DEVICES];
+	cl_int eventStatus = CL_QUEUED;
+	size_t globalThreads[1];
+	size_t localThreads[1];
+	size_t kernelWorkGroupSize;
+	cl_kernel kernel;
+
+	globalThreads[0] = ((config_opencl->width) * (config_opencl->height))
+			/ num_devices;
+	localThreads[0] = 256;
+
+	globalThreads[0] >>= 2;
+
+	for (cl_uint i = 0; i < num_devices; i++)
+	{
+		kernel = kernel_vector[i];
+		ret = clGetKernelWorkGroupInfo(kernel, devices[i],
+				CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &kernelWorkGroupSize,
+				0);
+		if ((cl_uint) (localThreads[0]) > kernelWorkGroupSize)
+		{
+			localThreads[0] = kernelWorkGroupSize;
+		}
+
+		xpos = 0.0;
+		ypos = 0.0;
+		xsize = 0.1;
+		ysize = 0.1;
+
+		xstep = (xsize / (double) config_opencl->width);
+		ystep = (ysize / (double) config_opencl->height);
+		leftx = (xpos - xsize / 2.0);
+		topy =
+				(ypos + ysize / 2.0
+						- ((double) i * ysize) / (double) num_devices);
+
+		if (i == 0)
+		{
+			topy0 = topy;
+		}
+
+		printf("xsize: %f, ysize: %f, xstep: %f, ystep: %f, leftx: %f, topy: %f\n", xsize, ysize, xstep, ystep, leftx, topy);
+
+		ret = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *) &outputBuffer[i]);
+
+		cl_float leftxF = (float) leftx;
+		cl_float topyF = (float) topy;
+		cl_float xstepF = (float) xstep;
+		cl_float ystepF = (float) ystep;
+
+		switch (config_opencl->fpu)
+		{
+		case OPENCL_FPU_32: // lel
+			ret = clSetKernelArg(kernel, 1, sizeof(cl_float), (void *) &leftxF);
+			ret = clSetKernelArg(kernel, 2, sizeof(cl_float), (void *) &topyF);
+			ret = clSetKernelArg(kernel, 3, sizeof(cl_float), (void *) &xstepF);
+			ret = clSetKernelArg(kernel, 4, sizeof(cl_float), (void *) &ystepF);
+			break;
+		case OPENCL_FPU_64:
+			ret = clSetKernelArg(kernel, 1, sizeof(cl_double), (void *) &leftx);
+			ret = clSetKernelArg(kernel, 2, sizeof(cl_double), (void *) &topy);
+			ret = clSetKernelArg(kernel, 3, sizeof(cl_double), (void *) &xstep);
+			ret = clSetKernelArg(kernel, 4, sizeof(cl_double), (void *) &ystep);
+			break;
+		case OPENCL_FPU_128:
+			break;
+		}
+
+		ret = clSetKernelArg(kernel, 5, sizeof(cl_uint),
+				(void *) &config_opencl->iterations);
+		ret = clSetKernelArg(kernel, 6, sizeof(cl_int),
+				(void *) &config_opencl->width);
+		ret = clEnqueueNDRangeKernel(commandQueue[i], kernel, 1, NULL,
+				globalThreads, localThreads, 0, NULL, &events[i]);
+
+
+		for (cl_uint i = 0; i < num_devices; i++)
+		{
+			ret = clFlush(commandQueue[i]);
+		}
+		for (cl_uint i = 0; i < num_devices; i++)
+		{
+			ret = clWaitForEvents(1, &events[num_devices - i - 1]);
+			ret = clReleaseEvent(events[num_devices - i - 1]);
+		}
+
+		for (cl_uint i = 0; i < num_devices; i++)
+		{
+			ret = clEnqueueReadBuffer(commandQueue[i], outputBuffer[i],
+			CL_FALSE, 0,
+					((config_opencl->width) * (config_opencl->height)
+							* sizeof(u32)) / num_devices,
+					config_opencl->arr
+							+ (((config_opencl->width) * (config_opencl->height) * i)
+									/ num_devices), 0,
+					NULL, &events[i]);
+		}
+
+		for (cl_uint i = 0; i < num_devices; i++)
+		{
+			ret = clFlush(commandQueue[i]);
+		}
+
+		for (cl_uint i = 0; i < num_devices; i++)
+		{
+			ret = clWaitForEvents(1, &events[num_devices - i - 1]);
+			ret = clReleaseEvent(events[num_devices - i - 1]);
+		}
+	}
+
+	glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT);
+	glBindTexture(GL_TEXTURE_2D, config_opencl->tex);
+	glEnable(GL_TEXTURE_2D);
+
+	glBegin(GL_QUADS);
+		glTexCoord2i(0, 0); glVertex2i(0, 0);
+		glTexCoord2i(0, 1); glVertex2i(0, config_opencl->height);
+		glTexCoord2i(1, 1); glVertex2i(config_opencl->width, config_opencl->height);
+		glTexCoord2i(1, 0); glVertex2i(config_opencl->width, 0);
+	glEnd();
 
+	glDisable(GL_TEXTURE_2D);
+	glBindTexture(GL_TEXTURE_2D, 0);
+	glutSwapBuffers();
 }
 
 void idle_opencl(void)
 {
+	static int t_old;
+	int t = 0, delta = 0;
+	do
+	{
+		t = glutGet(GLUT_ELAPSED_TIME);
+		delta = t - t_old;
+	} while (delta < 16); // TODO: Hardcoded FPS
+	t_old = t;
 
+	//glGenTextures(1, &tex);
+	glBindTexture(GL_TEXTURE_2D, config_opencl->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_opencl->width,
+			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);
+	glutPostRedisplay();
 }
diff --git a/src/render_opencl.h b/src/render_opencl.h
index aaa9df6..b995c4d 100644
--- a/src/render_opencl.h
+++ b/src/render_opencl.h
@@ -11,10 +11,24 @@
 #include "defs.h"
 #include <stdlib.h>
 #include <GL/glut.h>
+#include <CL/cl.h>
+#include "sets.h"
+#include <math.h>
+
+#define OPENCL_FPU_32 0
+#define OPENCL_FPU_64 1
+#define OPENCL_FPU_128 2
+
+#define MAX_SOURCE_SIZE 0xFFFF // 64 KiB
+#define MAX_DEVICES 4
 
 typedef struct config_opencl {
+	u8 fpu;
+	u8 fma;
+	GLuint tex;
+	d64 (*zoom_func)(d64, d64);
+	u8 set_func; // id, not pointer!
 	u32 *arr;
-	u32 (*set_func)(d64, d64, u32);
 	u32 iterations;
 	u32 colorFrom;
 	u32 colorTo;
@@ -28,6 +42,32 @@ 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;
+
+float cl_ft;
+cl_uint *output;
+cl_device_id device_id;
+cl_context context;
+cl_int ret;
+cl_kernel kernel_vector[MAX_DEVICES];
+cl_uint num_devices;
+cl_device_id *devices;
+cl_command_queue commandQueue[MAX_DEVICES];
+cl_mem outputBuffer[MAX_DEVICES];
+
+double xpos;
+double ypos;
+double xsize;
+double ysize;
+double xstep;
+double ystep;
+double leftx;
+double topy;
+double topy0;
+
+d64 zoom_func(d64 ft, d64 s);
+
 void init_opencl(OpenCLConfig *config);
 void render_opencl(void);
 void idle_opencl(void);
diff --git a/src/sets.h b/src/sets.h
index cea87fc..3dd8227 100644
--- a/src/sets.h
+++ b/src/sets.h
@@ -10,6 +10,14 @@
 
 #include "defs.h"
 
+#define SFUNC_MANDELBROT_R 0
+#define SFUNC_MANDELBROT_S (~0)
+#define SFUNC_JULIA_R 1
+#define SFUNC_JULIA_S (~1)
+
+#define SFUNC_MANDELBROT SFUNC_MANDELBROT_R
+#define SFUNC_JULIA SFUNC_JULIA_R
+
 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);