opencl example

CMakeLists.txt

@@ -1,22 +1,36 @@
 project(mandelbrot)
 
+set(CMAKE_MODULE_PATH "${CMAKE_SOURCE_DIR}/cmake")
+
+
 cmake_minimum_required(VERSION 2.8)
 find_package(PkgConfig REQUIRED)
 pkg_search_module(GLIB REQUIRED glib-2.0)
 pkg_check_modules(GTK3 REQUIRED gtk+-3.0)
 pkg_check_modules(CAIRO REQUIRED cairo)
 
+find_package(OpenCL REQUIRED)
+
+
+
 add_subdirectory(glade)
 
-include_directories(${GLIB_INCLUDE_DIRS} ${GTK3_INCLUDE_DIRS} ${CAIRO_INCLUDE_DIRS})
+include_directories(${GLIB_INCLUDE_DIRS} ${GTK3_INCLUDE_DIRS} ${CAIRO_INCLUDE_DIRS} ${OPENCL_INCLUDE_DIRS})
 link_directories(${GLIB_LINK_DIRS} ${GTK3_LINK_DIRS} ${CAIRO_LINK_DIRS})
 add_definitions(${GLIB2_CFLAGS_OTHER})
 
+configure_file("mandelbrot.cl" "cl_kernels/mandelbrot.cl" COPYONLY)
+
 aux_source_directory("src" SOURCES)
 add_compile_options(-Wall)
 
 add_executable(${PROJECT_NAME} ${SOURCES} ${CMAKE_CURRENT_BINARY_DIR}/glade/resources.c)
 add_dependencies(${PROJECT_NAME} glib-resources)
 SET_SOURCE_FILES_PROPERTIES(${CMAKE_CURRENT_BINARY_DIR}/glade/resources.c PROPERTIES GENERATED 1)
-target_link_libraries(${PROJECT_NAME} ${GLIB_LDFLAGS} ${GTK3_LDFLAGS} ${CAIRO_LDFLAGS} m pthread)
+target_link_libraries(${PROJECT_NAME} ${GLIB_LDFLAGS} ${GTK3_LDFLAGS} ${CAIRO_LDFLAGS} m pthread ${OPENCL_LIBRARIES})
 install (TARGETS ${PROJECT_NAME} DESTINATION bin)
+
+message(STATUS "OpenCL found: ${OPENCL_FOUND}")
+message(STATUS "OpenCL includes: ${OPENCL_INCLUDE_DIRS}")
+message(STATUS "OpenCL CXX includes: ${OPENCL_HAS_CPP_BINDINGS}")
+message(STATUS "OpenCL libraries: ${OPENCL_LIBRARIES}")

cmake/FindOpenCL.cmake

@@ -0,0 +1,171 @@
+# Find OpenCL
+#
+# To set manually the paths, define these environment variables:
+# OpenCL_INCPATH    - Include path (e.g. OpenCL_INCPATH=/opt/cuda/4.0/cuda/include)
+# OpenCL_LIBPATH    - Library path (e.h. OpenCL_LIBPATH=/usr/lib64/nvidia)
+#
+# Once done this will define
+#  OPENCL_FOUND            - system has OpenCL
+#  OPENCL_INCLUDE_DIRS     - the OpenCL include directory
+#  OPENCL_LIBRARIES        - link these to use OpenCL
+#  OPENCL_HAS_CPP_BINDINGS - system has also cl.hpp
+
+FIND_PACKAGE(PackageHandleStandardArgs)
+
+SET (OPENCL_VERSION_STRING "0.1.0")
+SET (OPENCL_VERSION_MAJOR 0)
+SET (OPENCL_VERSION_MINOR 1)
+SET (OPENCL_VERSION_PATCH 0)
+
+IF (APPLE)
+
+	# IF OpenCL_LIBPATH is given use it and don't use default path
+	IF (DEFINED ENV{OpenCL_LIBPATH})
+		FIND_LIBRARY(OPENCL_LIBRARIES OpenCL PATHS ENV OpenCL_LIBPATH NO_DEFAULT_PATH)
+	ELSE ()
+		FIND_LIBRARY(OPENCL_LIBRARIES OpenCL DOC "OpenCL lib for OSX")
+	ENDIF ()
+
+	# IF OpenCL_INCPATH is given use it and find for CL/cl.h and OpenCL/cl.h do not try to find default paths
+	IF (DEFINED ENV{OpenCL_INCPATH})
+		FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h OpenCL/cl.h PATHS ENV OpenCL_INCPATH NO_DEFAULT_PATH)
+		FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp OpenCL/cl.hpp PATHS ${OPENCL_INCLUDE_DIRS} NO_DEFAULT_PATH)
+	ELSE ()
+		FIND_PATH(OPENCL_INCLUDE_DIRS OpenCL/cl.h DOC "Include for OpenCL on OSX")
+		FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS OpenCL/cl.hpp DOC "Include for OpenCL CPP bindings on OSX")
+	ENDIF ()
+
+ELSE (APPLE)
+
+	IF (WIN32)
+
+		# Find OpenCL includes and libraries from environment variables provided by vendor
+		SET(OPENCL_INCLUDE_SEARCH_PATHS)
+		SET(OPENCL_LIBRARY_SEARCH_PATHS)
+		SET(OPENCL_LIBRARY_64_SEARCH_PATHS)
+
+		# Nvidia
+		IF (DEFINED ENV{CUDA_INC_PATH})
+			SET(OPENCL_INCLUDE_SEARCH_PATHS ${OPENCL_INCLUDE_SEARCH_PATHS} $ENV{CUDA_INC_PATH})
+			SET(OPENCL_LIBRARY_64_SEARCH_PATHS ${OPENCL_LIBRARY_64_SEARCH_PATHS} $ENV{CUDA_LIB_PATH}/../lib64)
+			SET(OPENCL_LIBRARY_SEARCH_PATHS ${OPENCL_LIBRARY_SEARCH_PATHS} $ENV{CUDA_LIB_PATH}/../lib)
+		ENDIF()
+		IF (DEFINED ENV{CUDA_PATH})
+			SET(OPENCL_INCLUDE_SEARCH_PATHS ${OPENCL_INCLUDE_SEARCH_PATHS} $ENV{CUDA_INC_PATH})
+			SET(OPENCL_LIBRARY_64_SEARCH_PATHS ${OPENCL_LIBRARY_64_SEARCH_PATHS} $ENV{CUDA_PATH}/lib/x64/)
+			SET(OPENCL_LIBRARY_SEARCH_PATHS ${OPENCL_LIBRARY_SEARCH_PATHS} $ENV{CUDA_PATH}/lib/Win32/)
+		ENDIF()
+
+		# Intel SDK
+		IF (DEFINED ENV{INTELOCSDKROOT})
+			SET(OPENCL_INCLUDE_SEARCH_PATHS ${OPENCL_INCLUDE_SEARCH_PATHS} $ENV{INTELOCSDKROOT}/include)
+			SET(OPENCL_LIBRARY_64_SEARCH_PATHS ${OPENCL_LIBRARY_64_SEARCH_PATHS} $ENV{INTELOCSDKROOT}/lib/x64)
+			SET(OPENCL_LIBRARY_SEARCH_PATHS ${OPENCL_LIBRARY_SEARCH_PATHS} $ENV{INTELOCSDKROOT}/lib/x86)
+		ENDIF()
+
+		# AMD SDK
+		IF (DEFINED ENV{AMDAPPSDKROOT})
+			SET(OPENCL_INCLUDE_SEARCH_PATHS ${OPENCL_INCLUDE_SEARCH_PATHS} $ENV{AMDAPPSDKROOT}/include)
+			SET(OPENCL_LIBRARY_64_SEARCH_PATHS ${OPENCL_LIBRARY_64_SEARCH_PATHS} $ENV{AMDAPPSDKROOT}/lib/x86_64)
+			SET(OPENCL_LIBRARY_SEARCH_PATHS ${OPENCL_LIBRARY_SEARCH_PATHS} $ENV{AMDAPPSDKROOT}/lib/x86)
+		ENDIF()
+
+		# Override search paths with OpenCL_INCPATH env variable
+		IF (DEFINED ENV{OpenCL_INCPATH})
+			SET(OPENCL_INCLUDE_SEARCH_PATHS $ENV{OpenCL_INCPATH})
+		ENDIF ()
+
+		# Override search paths with OpenCL_LIBPATH env variable
+		IF (DEFINED ENV{OpenCL_LIBPATH})
+			SET(OPENCL_LIBRARY_SEARCH_PATHS $ENV{OpenCL_LIBPATH})
+			SET(OPENCL_LIBRARY_64_SEARCH_PATHS $ENV{OpenCL_LIBPATH})
+		ENDIF ()
+
+		FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h PATHS ${OPENCL_INCLUDE_SEARCH_PATHS})
+		FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp PATHS ${OPENCL_INCLUDE_SEARCH_PATHS})
+		
+		FIND_LIBRARY(_OPENCL_32_LIBRARIES OpenCL.lib HINTS ${OPENCL_LIBRARY_SEARCH_PATHS} PATHS ${OPENCL_LIB_DIR} ENV PATH)
+		FIND_LIBRARY(_OPENCL_64_LIBRARIES OpenCL.lib HINTS ${OPENCL_LIBRARY_64_SEARCH_PATHS} PATHS ${OPENCL_LIB_DIR} ENV PATH)
+
+		# Check if 64bit or 32bit versions links fine
+		SET (_OPENCL_VERSION_SOURCE "${CMAKE_BINARY_DIR}${CMAKE_FILES_DIRECTORY}/openclversion.c")
+		#SET (_OPENCL_VERSION_SOURCE "${CMAKE_BINARY_DIR}/test.c")
+		FILE (WRITE "${_OPENCL_VERSION_SOURCE}"
+			"
+			#if __APPLE__
+			#include <OpenCL/cl.h>
+			#else /* !__APPLE__ */
+			#include <CL/cl.h>
+			#endif /* __APPLE__ */
+			int main()
+			{	
+			    cl_int result;
+			    cl_platform_id id;
+			    result = clGetPlatformIDs(1, &id, NULL);
+			    return result != CL_SUCCESS;
+			}
+			")
+
+		TRY_COMPILE(_OPENCL_64_COMPILE_SUCCESS ${CMAKE_BINARY_DIR} "${_OPENCL_VERSION_SOURCE}"
+			CMAKE_FLAGS
+			"-DINCLUDE_DIRECTORIES:STRING=${OPENCL_INCLUDE_DIRS}"
+			CMAKE_FLAGS
+			"-DLINK_LIBRARIES:STRING=${_OPENCL_64_LIBRARIES}"
+		)
+
+		IF(_OPENCL_64_COMPILE_SUCCESS)
+			message(STATUS "OpenCL 64bit lib found.")
+			SET(OPENCL_LIBRARIES ${_OPENCL_64_LIBRARIES})
+		ELSE()
+	  		TRY_COMPILE(_OPENCL_32_COMPILE_SUCCESS ${CMAKE_BINARY_DIR} "${_OPENCL_VERSION_SOURCE}"
+				CMAKE_FLAGS
+				"-DINCLUDE_DIRECTORIES:STRING=${OPENCL_INCLUDE_DIRS}"
+				CMAKE_FLAGS
+				"-DLINK_LIBRARIES:STRING=${_OPENCL_32_LIBRARIES}"
+	  		)
+			IF(_OPENCL_32_COMPILE_SUCCESS)
+				message(STATUS "OpenCL 32bit lib found.")
+				SET(OPENCL_LIBRARIES ${_OPENCL_32_LIBRARIES})
+			ELSE()
+				message(STATUS "Couldn't link opencl..")
+			ENDIF()
+		ENDIF()
+
+
+	ELSE (WIN32)
+
+		IF (CYGWIN)
+			SET (CMAKE_FIND_LIBRARY_SUFFIXES .lib)
+			SET (OCL_LIB_SUFFIX .lib)
+		ENDIF (CYGWIN)
+
+		# Unix style platforms
+		FIND_LIBRARY(OPENCL_LIBRARIES OpenCL${OCL_LIB_SUFFIX}
+			PATHS ENV LD_LIBRARY_PATH ENV OpenCL_LIBPATH
+		)
+
+		GET_FILENAME_COMPONENT(OPENCL_LIB_DIR ${OPENCL_LIBRARIES} PATH)
+		GET_FILENAME_COMPONENT(_OPENCL_INC_CAND ${OPENCL_LIB_DIR}/../../include ABSOLUTE)
+
+		# The AMD SDK currently does not place its headers
+		# in /usr/include, therefore also search relative
+		# to the library
+		FIND_PATH(OPENCL_INCLUDE_DIRS CL/cl.h PATHS ${_OPENCL_INC_CAND} "/usr/local/cuda/include" "/opt/AMDAPP/include" ENV OpenCL_INCPATH)
+		FIND_PATH(_OPENCL_CPP_INCLUDE_DIRS CL/cl.hpp PATHS ${_OPENCL_INC_CAND} "/usr/local/cuda/include" "/opt/AMDAPP/include" ENV OpenCL_INCPATH)
+
+	ENDIF (WIN32)
+
+ENDIF (APPLE)
+
+FIND_PACKAGE_HANDLE_STANDARD_ARGS(OpenCL DEFAULT_MSG OPENCL_LIBRARIES OPENCL_INCLUDE_DIRS)
+
+IF(_OPENCL_CPP_INCLUDE_DIRS)
+	SET( OPENCL_HAS_CPP_BINDINGS TRUE )
+	LIST( APPEND OPENCL_INCLUDE_DIRS ${_OPENCL_CPP_INCLUDE_DIRS} )
+	# This is often the same, so clean up
+	LIST( REMOVE_DUPLICATES OPENCL_INCLUDE_DIRS )
+ENDIF(_OPENCL_CPP_INCLUDE_DIRS)
+
+MARK_AS_ADVANCED(
+  OPENCL_INCLUDE_DIRS
+)

mandelbrot.cl

@@ -0,0 +1,34 @@
+__kernel void mandelbrot(__global unsigned int *out_buff, unsigned int iter, double cx, double cy, double xs, double ys, unsigned int width, unsigned int height)
+{
+	const int idx = get_global_id(0);
+
+	unsigned int i;
+	double z_real, z_imag;
+	double z2_real, z2_imag;
+	double pt_real, pt_imag;
+
+	const unsigned int pix_x = idx % width;
+	const unsigned int pix_y = (unsigned int)idx / width;
+	const double val_per_x_pixel = xs / (width - 1);
+	const double val_per_y_pixel = ys / (height - 1); 
+	
+
+	pt_real = (((double)pix_x - (double)(width - 1) / 2)) * val_per_x_pixel + cx;
+	pt_imag = (((double)pix_y - (double)(height - 1) / 2)) * val_per_y_pixel + cy;
+
+	z_real = 0.0;
+        z_imag = 0.0;	
+
+	for (i = 0; i < iter; i++) {
+		z2_real = z_real * z_real - z_imag * z_imag;
+		z2_imag = 2.0 * z_real * z_imag;
+
+		z_real = z2_real + pt_real;
+		z_imag = z2_imag + pt_imag;
+
+		if ((z_real * z_real + z_imag * z_imag) >= 4.0)
+			break;
+	}
+
+	out_buff[idx] = i;
+}

src/main.c

@@ -1,17 +1,19 @@
 #include <stdio.h>
 #include <gtk/gtk.h>
 #include <math.h>
-#include <threads.h>
+#include <errno.h>
+#include <stdlib.h>
+#include <CL/cl.h>
 
 struct canvas_buffer {
-	unsigned int *mandelbrot_buffer;
-	unsigned int iterations;
-	unsigned int height;
-	unsigned int width;
-	double center_x;
-	double center_y;
-	double x_span;
-	double y_span;
+		unsigned int *mandelbrot_buffer;
+		unsigned int iterations;
+		unsigned int height;
+		unsigned int width;
+		double center_x;
+		double center_y;
+		double x_span;
+		double y_span;
 };
 
 static gboolean on_main_window_close(GtkWidget *window, gpointer data)
@@ -20,6 +22,43 @@ static gboolean on_main_window_close(GtkWidget *window, gpointer data)
 	return FALSE;
 }
 
+void check_error(cl_int error)
+{
+	if (error != CL_SUCCESS) {
+		printf("OpenCL call failed with error %d\n", (int)error);
+		exit(error);
+	}
+}
+
+void load_kernel_from_file(char *file, char **src)
+{
+	long fsize;
+	FILE *f = fopen(file, "rb");
+
+	fseek(f, 0, SEEK_END);
+	fsize = ftell(f);
+	fseek(f, 0, SEEK_SET);  //same as rewind(f);
+	*src = (char *)malloc(fsize + 1);
+	fread(*src, fsize, 1, f);
+	fclose(f);
+	(*src)[fsize] = 0;
+}
+
+cl_program create_program(char *src, cl_context context)
+{
+	cl_program program;
+	size_t lengths[1];
+	cl_int error = 0;
+
+	lengths[0] = strlen(src);
+	printf("strlen: %u\n", (unsigned int)lengths[0]);
+	program = clCreateProgramWithSource(context, 1, (const char **)&src,
+					    lengths, &error);
+	check_error(error);
+
+	return program;
+}
+unsigned int *data_array = NULL;
 static gboolean drawing_callback(GtkWidget *widget, cairo_t *cr, gpointer data)
 {
 	struct canvas_buffer *buff = (struct canvas_buffer *)data;
@@ -38,24 +77,24 @@ static gboolean drawing_callback(GtkWidget *widget, cairo_t *cr, gpointer data)
 	cairo_scale(cr, 1, -1);
 	cairo_translate(cr, (double)width/2.0, -(double)height/2.0);
 
-	for (x = 0; x < buff->width; x++) {
-		for (y = 0; y < buff->height; y++) {
-			cairo_rectangle(cr,
-					((double)x - (double)(buff->width - 1) / 2),
-					((double)y - (double)(buff->height - 1) / 2),
-					1,
-					1);
-			if (buff->mandelbrot_buffer[y * buff->width + x] == buff->iterations) {
-				cairo_set_source_rgb (cr, 0, 0, 0);
-			} else {
-				color_scale = (double)buff->mandelbrot_buffer[y * buff->width + x] / buff->iterations;
-				cairo_set_source_rgb (cr, (sqrt(1-color_scale)), (color_scale*color_scale), color_scale);
-			}
-			cairo_fill(cr);
 
-		}
-	}
+       for (x = 0; x < buff->width; x++) {
+	       for (y = 0; y < buff->height; y++) {
+		       cairo_rectangle(cr,
+				       ((double)x - (double)(buff->width - 1) / 2),
+				       ((double)y - (double)(buff->height - 1) / 2),
+				      1,
+				       1);
+		       if (data_array[y * buff->width + x] == buff->iterations) {
+			       cairo_set_source_rgb (cr, 0, 0, 0);
+		      } else {
+			       color_scale = (double)data_array[y * buff->width + x] / buff->iterations;
+			       cairo_set_source_rgb (cr, (sqrt(1-color_scale)), (color_scale*color_scale), color_scale);
+		       }
+		       cairo_fill(cr);
 
+	       }
+	       }
 
 
 	/* Drawing code end */
@@ -64,58 +103,124 @@ static gboolean drawing_callback(GtkWidget *widget, cairo_t *cr, gpointer data)
 	return FALSE;
 }
 
-static void calculate_mandelbrot(struct canvas_buffer *buff)
+
+static int calculate_mandelbrot(struct canvas_buffer *buff)
 {
-	double val_per_x_pixel;
-	double val_per_y_pixel;
-	unsigned int i;
-	unsigned int pixel_x;
-	unsigned int pixel_y;
-	double point_real, point_imag;
-	double z_real, z_imag;
-	double z2_real, z2_imag;
+
+	int i;
+	cl_int error;
+	cl_int ret;
+	cl_platform_id *platform_ids = NULL;
+	cl_device_id *device_ids = NULL;
+	cl_uint platform_id_count = 0;
+	cl_uint device_id_count = 0;
+	cl_context context;
+	cl_program program;
+	cl_kernel kernel;
+	cl_command_queue queue;
+	cl_mem dat_buff;
+	char *source_code = NULL;
+	char temp_buff[100];
+	size_t data_size;
+	size_t work_size = 64;
 
 
-	if (!buff)
-		return;
+	data_size = buff->width*buff->height;
+	data_array = (unsigned int *)malloc(sizeof(data_array[0])*data_size);
+	if (!data_array)
+		return -ENOMEM;
 
-	/* Setup the calculation data */
-	val_per_x_pixel = (buff->x_span) / (buff->width - 1);
-	val_per_y_pixel = (buff->y_span) / (buff->height - 1);
 
-	for (pixel_y = 0; pixel_y < buff->height; pixel_y++) {
-		for (pixel_x = 0; pixel_x < buff->width; pixel_x++) {
-			point_real = (((double)pixel_x - (double)(buff->width - 1) / 2)) * val_per_x_pixel + buff->center_x;
-			point_imag = (((double)pixel_y - (double)(buff->height - 1) / 2)) * val_per_y_pixel + buff->center_y;
+	clGetPlatformIDs(0, NULL, &platform_id_count);
+	if (platform_id_count == 0)
+		return -ENOMEM;
 
-			z_real = 0.0;
-			z_imag = 0.0;
+	platform_ids = (cl_platform_id *)malloc(platform_id_count *
+						sizeof(cl_platform_id));
+	if (platform_ids == NULL)
+		return -ENOMEM;
 
-			/* Check for convergence */
-			for (i = 0; i < buff->iterations; i++) {
+	clGetPlatformIDs(platform_id_count, platform_ids, NULL);
 
-				z2_real = z_real * z_real - z_imag * z_imag;
-				z2_imag = 2.0 * z_real * z_imag;
+	printf("Platforms available: %u\n", (unsigned int)platform_id_count);
 
-				z_real = z2_real + point_real;
-				z_imag = z2_imag + point_imag;
+	for (i = 0; i < platform_id_count; i++) {
+		temp_buff[0] = '\0';
+		clGetPlatformInfo(platform_ids[i], CL_PLATFORM_NAME, sizeof(temp_buff), temp_buff, NULL);
+		printf("Platform ID %d: %s\n", i, temp_buff);
+	}
 
-				if ((z_real * z_real + z_imag * z_imag) >= 4.0)
-					break;
-			}
+	clGetDeviceIDs(platform_ids[0], CL_DEVICE_TYPE_ALL, 0, NULL,
+			&device_id_count);
+	printf("Device in Platform 0 count: %u\n",
+	       (unsigned int)device_id_count);
 
-			buff->mandelbrot_buffer[pixel_y * buff->width + pixel_x] = i;
-		}
+	if (device_id_count == 0)
+		return -1;
+
+	device_ids = (cl_device_id *) malloc(sizeof(cl_device_id)
+					     *device_id_count);
+
+	clGetDeviceIDs(platform_ids[0], CL_DEVICE_TYPE_ALL, device_id_count,
+			device_ids, NULL);
+
+	for (i = 0; i < device_id_count; i++) {
+		temp_buff[0] = '\0';
+		clGetDeviceInfo(device_ids[i], CL_DEVICE_NAME, sizeof(temp_buff), temp_buff, NULL);
+		printf("Device ID %d: %s\n", i, temp_buff);
 	}
 
 
+	context = clCreateContext(NULL, 1, &device_ids[0], NULL, NULL, &error);
+	check_error(error);
+
+	queue = clCreateCommandQueue(context, device_ids[0], 0, &error);
+
+	dat_buff = clCreateBuffer(context, CL_MEM_WRITE_ONLY |
+					CL_MEM_COPY_HOST_PTR,
+					sizeof(data_array[0])*data_size, data_array, &error);
+	check_error(error);
+
+	load_kernel_from_file("cl_kernels/mandelbrot.cl", &source_code);
+	program = create_program(source_code, context);
+
+	/* Compile the source code */
+	check_error(clBuildProgram(program, device_id_count,
+					   device_ids, NULL, NULL, NULL));
+	kernel = clCreateKernel(program, "mandelbrot", &error);
+
+
+	/* Setting the arguments */
+	clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&dat_buff);
+	clSetKernelArg(kernel, 1, sizeof(unsigned int), (void *)&buff->iterations);
+	clSetKernelArg(kernel, 2, sizeof(double), (void *)&buff->center_x);
+	clSetKernelArg(kernel, 3, sizeof(double), (void *)&buff->center_y);
+	clSetKernelArg(kernel, 4, sizeof(double), (void *)&buff->x_span);
+	clSetKernelArg(kernel, 5, sizeof(double), (void *)&buff->y_span);
+	clSetKernelArg(kernel, 6, sizeof(unsigned int), (void *)&buff->width);
+	clSetKernelArg(kernel, 7, sizeof(unsigned int), (void *)&buff->height);
+
+	ret = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &data_size, &work_size, 0, NULL, NULL);
+		check_error(ret);
+
+	clEnqueueReadBuffer(queue, dat_buff, CL_TRUE, 0, data_size*sizeof(data_array[0]), data_array, 0, NULL, NULL);
+
+	ret = clFlush(queue);
+	ret = clFinish(queue);
+	ret = clReleaseKernel(kernel);
+	ret = clReleaseProgram(program);
+	ret = clReleaseMemObject(dat_buff);
+	ret = clReleaseCommandQueue(queue);
+	ret = clReleaseContext(context);
+	free(platform_ids);
+	free(device_ids);
+
+	return 0;
+
 }
 
 
-int f(void * a)
-{
-	return 1;
-}
+
 
 int main(int argc, char **argv)
 {
@@ -133,22 +238,19 @@ int main(int argc, char **argv)
 
 	g_signal_connect(window, "delete-event", G_CALLBACK(on_main_window_close), NULL);
 	g_signal_connect(G_OBJECT(drawing_area),
-			"draw", G_CALLBACK(drawing_callback), (gpointer)&mandelbrot_buff);
+			 "draw", G_CALLBACK(drawing_callback), (gpointer)&mandelbrot_buff);
 
 	g_object_unref(builder);
 
-	 thrd_t t;
-	thrd_create(&t, f, NULL);
 
-
-	mandelbrot_buff.x_span = 2.8;
-	mandelbrot_buff.y_span = 2.25;
-	mandelbrot_buff.center_x = -0.5;
-	mandelbrot_buff.center_y = 0.0;
-	mandelbrot_buff.width = 1900;
-	mandelbrot_buff.height = 1000;
-	mandelbrot_buff.iterations = 300;
-	mandelbrot_buff.mandelbrot_buffer = (unsigned int *)malloc(sizeof(unsigned int) * mandelbrot_buff.height * mandelbrot_buff.width);
+	mandelbrot_buff.x_span = 1;//2.8;
+	mandelbrot_buff.y_span = 0.5;//2.25;
+	mandelbrot_buff.center_x = -0.6;
+	mandelbrot_buff.center_y = 0.6;
+	mandelbrot_buff.width = 4000;
+	mandelbrot_buff.height = 2000;
+	mandelbrot_buff.iterations = 70;
+	mandelbrot_buff.mandelbrot_buffer = NULL;
 	calculate_mandelbrot(&mandelbrot_buff);