From c8a38bd8ad66eb0b35f08a4733fdee37a888b83c Mon Sep 17 00:00:00 2001 From: Leonard Kugis Date: Sat, 27 Jan 2018 02:03:31 +0100 Subject: OpenCL rendering added, not working properly --- src/cl/mandelbrot32.cl | 196 ++++++++++++++++++++++++++++++++++ src/cl/mandelbrot64.cl | 230 ++++++++++++++++++++++++++++++++++++++++ src/defs.h | 2 + src/mandelbrot-zoom.c | 95 ++++++++--------- src/mandelbrot-zoom.h | 3 + src/mandelbrot.cl | 6 -- src/render.c | 34 ++++-- src/render.h | 6 +- src/render_cpu.c | 36 ++++--- src/render_cpu.h | 9 +- src/render_opencl.c | 277 ++++++++++++++++++++++++++++++++++++++++++++++++- src/render_opencl.h | 42 +++++++- src/sets.h | 8 ++ 13 files changed, 851 insertions(+), 93 deletions(-) create mode 100644 src/cl/mandelbrot32.cl create mode 100644 src/cl/mandelbrot64.cl delete mode 100644 src/mandelbrot.cl (limited to 'src') 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 + #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 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 #include #include +#include #include #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 #include +#include +#include "sets.h" +#include + +#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); -- cgit v1.2.1