Merge branch 'master' of https://collaborating.tuhh.de/cev7691/mandelbrot-zoom

1 files changed, 297 insertions, 0 deletions

diff --git a/src/render_opencl.c b/src/render_opencl.c
new file mode 100644
index 0000000..9d650c2
--- /dev/null
+++ b/src/render_opencl.c
@@ -0,0 +1,297 @@
+/*
+ * render_opencl.c
+ *
+ *  Created on: 26.01.2018
+ *      Author: Superleo1810
+ */
+
+#include "render_opencl.h"
+
+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;
+	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();
+}