1
0
Prechádzať zdrojové kódy

DoubleFloat working

Nicolas Winkler 5 rokov pred
rodič
commit
b1520ca968

+ 1 - 40
choosegenerators.cpp

@@ -97,7 +97,7 @@ double Benchmarker::benchmarkResult(mnd::Generator& mg) const
             break;
         }
         else if (time < std::chrono::milliseconds(10)) {
-            i += 3;
+            i += 7;
         }
     }
 
@@ -144,45 +144,6 @@ ChooseGenerators::ChooseGenerators(mnd::MandelContext& mndCtxt, QWidget *parent)
         }
     }
 
-    /*generators = std::map<QString, mnd::Generator*> {
-        { "float", mndCtxt.getCpuGenerator(mnd::GeneratorType::FLOAT) },
-        { "double", mndCtxt.getCpuGenerator(mnd::GeneratorType::DOUBLE) },
-        { "double double", mndCtxt.getCpuGenerator(mnd::GeneratorType::DOUBLE_DOUBLE) },
-        { "quad double", mndCtxt.getCpuGenerator(mnd::GeneratorType::QUAD_DOUBLE) },
-        { "float256", mndCtxt.getCpuGenerator(mnd::GeneratorType::FLOAT256) },
-        { "fixed512", mndCtxt.getCpuGenerator(mnd::GeneratorType::FIXED512) },
-    };
-
-    if (mndCtxt.getCpuInfo().hasSse2()) {
-        generators.insert({ "float SSE2", mndCtxt.getCpuGenerator(mnd::GeneratorType::FLOAT_SSE2) });
-        generators.insert({ "double SSE2", mndCtxt.getCpuGenerator(mnd::GeneratorType::DOUBLE_SSE2) });
-    }
-    if (mndCtxt.getCpuInfo().hasAvx()) {
-        generators.insert({ "float AVX", mndCtxt.getCpuGenerator(mnd::GeneratorType::FLOAT_AVX) });
-        generators.insert({ "double AVX", mndCtxt.getCpuGenerator(mnd::GeneratorType::DOUBLE_AVX) });
-        if (mndCtxt.getCpuInfo().hasFma()) {
-            generators.insert({ "double double AVX", mndCtxt.getCpuGenerator(mnd::GeneratorType::DOUBLE_DOUBLE_AVX) });
-        }
-    }
-    if (mndCtxt.getCpuInfo().hasNeon()) {
-        generators.insert({ "float Neon", mndCtxt.getCpuGenerator(mnd::GeneratorType::FLOAT_NEON) });
-        generators.insert({ "double Neon", mndCtxt.getCpuGenerator(mnd::GeneratorType::DOUBLE_NEON) });
-    }
-    for (auto& device : mndCtxt.getDevices()) {
-        if (mnd::Generator* gen; (gen = device.getGenerator(mnd::GeneratorType::FLOAT))) {
-            generators.insert({ QString("float ") + QString::fromStdString(device.getName()),
-                                gen });
-        }
-        if (mnd::Generator* gen; (gen = device.getGenerator(mnd::GeneratorType::DOUBLE))) {
-            generators.insert({ QString("double ") + QString::fromStdString(device.getName()),
-                                gen });
-        }
-        if (mnd::Generator* gen; (gen = device.getGenerator(mnd::GeneratorType::DOUBLE_DOUBLE))) {
-            generators.insert({ QString("double double ") + QString::fromStdString(device.getName()),
-                                gen });
-        }
-    }*/
-
     auto& defGen = mndCtxt.getDefaultGenerator();
     for (auto it = defGen.getGenerators().rbegin(); it != defGen.getGenerators().rend(); it++) {
         auto& [prec, gen] = *it;

+ 3 - 3
choosegenerators.ui

@@ -34,15 +34,15 @@
            <property name="alternatingRowColors">
             <bool>true</bool>
            </property>
+           <attribute name="horizontalHeaderVisible">
+            <bool>true</bool>
+           </attribute>
            <attribute name="horizontalHeaderDefaultSectionSize">
             <number>180</number>
            </attribute>
            <attribute name="horizontalHeaderMinimumSectionSize">
             <number>100</number>
            </attribute>
-           <attribute name="horizontalHeaderStretchLastSection">
-            <bool>false</bool>
-           </attribute>
            <attribute name="verticalHeaderVisible">
             <bool>false</bool>
            </attribute>

+ 2 - 0
libmandel/include/OpenClCode.h

@@ -3,6 +3,8 @@
 
 namespace mnd
 {
+    const char* getFloat_cl();
+    const char* getDoubleFloat_cl();
     const char* getFixed64_cl();
     const char* getFixed128_cl();
     const char* getFixed512_cl();

+ 12 - 17
libmandel/src/ClGenerators.cpp

@@ -1,6 +1,5 @@
 #include "ClGenerators.h"
 #include "doubledouble.h"
-#include "doublefloat.h"
 #include "OpenClCode.h"
 
 #ifdef WITH_OPENCL
@@ -99,7 +98,9 @@ void ClGenerator::generate(const mnd::MandelInfo& info, float* data)
     float pixelScaleX = float(info.view.width / info.bWidth);
     float pixelScaleY = float(info.view.height / info.bHeight);
 
-    Kernel iterate = Kernel(program, "iterate");
+    bool useVec = device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT>() >= 4;
+
+    Kernel iterate = Kernel(program, useVec ? "iterate_vec4" : "iterate");
     iterate.setArg(0, buffer_A);
     iterate.setArg(1, int(info.bWidth));
     iterate.setArg(2, float(info.view.x));
@@ -109,8 +110,7 @@ void ClGenerator::generate(const mnd::MandelInfo& info, float* data)
     iterate.setArg(6, int(info.maxIter));
     iterate.setArg(7, int(info.smooth ? 1 : 0));
 
-    // TODO check for overflow
-    if (true || device.getInfo<CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT>() >= 4) {
+    if (useVec) {
         queue.enqueueNDRangeKernel(iterate, 0, NDRange(info.bWidth * info.bHeight / 4));
     } else {
         queue.enqueueNDRangeKernel(iterate, 0, NDRange(info.bWidth * info.bHeight));
@@ -140,7 +140,8 @@ ClGeneratorFloat::ClGeneratorFloat(cl::Device device) :
 
 std::string ClGeneratorFloat::getKernelCode(bool smooth) const
 {
-    return 
+    return mnd::getFloat_cl();
+        /*
 //        "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
         "__kernel void iterate(__global float* A, const int width, float xl, float yt, float pixelScaleX, float pixelScaleY, int max, int smooth) {"
         "   int index = get_global_id(0) * 4;\n"
@@ -177,7 +178,7 @@ std::string ClGeneratorFloat::getKernelCode(bool smooth) const
         "      else\n"
         "          A[index + i] = ((float) count[i]);\n"
         "   }"
-        "}";
+        "}";*/
         /*
 //        "#pragma OPENCL EXTENSION cl_khr_fp64 : enable"
         "__kernel void iterate(__global float* A, const int width, float xl, float yt, float pixelScaleX, float pixelScaleY, int max, int smooth) {"
@@ -271,7 +272,7 @@ std::pair<float, float> add(std::pair<float, float> a, std::pair<float, float> b
 
 std::pair<float, float> mul(std::pair<float, float> a, std::pair<float, float> b) {
     auto t = twoProd(a.first, b.first);
-    float t3 = ((a.first * b.second) + (a.second * b.first)) + t.second;
+    t.second += ((a.first * b.second) + (a.second * b.first));
     return twoSum(t.first, t.second);
 }
 
@@ -286,13 +287,6 @@ void ClGeneratorDoubleFloat::generate(const mnd::MandelInfo& info, float* data)
 {
     ::size_t bufferSize = info.bWidth * info.bHeight * sizeof(float);
 
-    auto add12 = [](float a, float b) {
-        float s = a + b;
-        float v = s - a;
-        float r = (a - (s - v)) + (b - v);
-        return std::pair{ s, r };
-    };
-
     auto splitDouble = [] (double x) {
         /*uint64_t xl = *((uint64_t*)&x);
         uint64_t mantissa = xl & 0x000FFFFFFFFFFFFFULL;
@@ -307,7 +301,7 @@ void ClGeneratorDoubleFloat::generate(const mnd::MandelInfo& info, float* data)
             //printf("hi: %.10ef, lo: %.10ef\n", hi, lo);
             //fflush(stdout);
         }
-        return std::pair{ hi, 0.0f };
+        return std::pair{ hi, lo };
     };
 
     Buffer buffer_A(context, CL_MEM_WRITE_ONLY, bufferSize);
@@ -319,6 +313,7 @@ void ClGeneratorDoubleFloat::generate(const mnd::MandelInfo& info, float* data)
     auto[w1, w2] = splitDouble(pixelScX);
     auto[h1, h2] = splitDouble(pixelScY);
 
+    /*
     for (int px = 0; px < info.bWidth; px++) {
         for (int py = 0; py < info.bHeight; py++) {
             std::pair<float, float> xl = { x1, x2 };
@@ -355,7 +350,7 @@ void ClGeneratorDoubleFloat::generate(const mnd::MandelInfo& info, float* data)
         }
     }
     return;
-   
+    */
     
     Kernel iterate = Kernel(program, "iterate");
     iterate.setArg(0, buffer_A);
@@ -378,7 +373,7 @@ void ClGeneratorDoubleFloat::generate(const mnd::MandelInfo& info, float* data)
 
 std::string ClGeneratorDoubleFloat::getKernelCode(bool smooth) const
 {
-    return (char*) doublefloat_cl;
+    return getDoubleFloat_cl();
 }
 
 

+ 55 - 25
libmandel/src/CpuGeneratorsAVX.cpp

@@ -52,25 +52,40 @@ void CpuGenerator<float, mnd::X86_AVX, parallel>::generate(const mnd::MandelInfo
 
             __m256 a = xs;
             __m256 b = ys;
-
-            for (int k = 0; k < info.maxIter; k++) {
-                __m256 aa = _mm256_mul_ps(a, a);
-                __m256 bb = _mm256_mul_ps(b, b);
-                __m256 abab = _mm256_mul_ps(a, b); abab = _mm256_add_ps(abab, abab);
-                a = _mm256_add_ps(_mm256_sub_ps(aa, bb), xs);
-                b = _mm256_add_ps(abab, ys);
-                __m256 cmp = _mm256_cmp_ps(_mm256_add_ps(aa, bb), threshold, _CMP_LE_OQ);
-                if (info.smooth) {
+            if (info.smooth) {
+                for (int k = 0; k < info.maxIter; k++) {
+                    __m256 aa = _mm256_mul_ps(a, a);
+                    __m256 bb = _mm256_mul_ps(b, b);
+                    __m256 abab = _mm256_mul_ps(a, b); abab = _mm256_add_ps(abab, abab);
+                    a = _mm256_add_ps(_mm256_sub_ps(aa, bb), xs);
+                    b = _mm256_add_ps(abab, ys);
+                    __m256 cmp = _mm256_cmp_ps(_mm256_add_ps(aa, bb), threshold, _CMP_LE_OQ);
                     resultsa = _mm256_or_ps(_mm256_andnot_ps(cmp, resultsa), _mm256_and_ps(cmp, a));
                     resultsb = _mm256_or_ps(_mm256_andnot_ps(cmp, resultsb), _mm256_and_ps(cmp, b));
+                    adder = _mm256_and_ps(adder, cmp);
+                    counter = _mm256_add_ps(counter, adder);
+                    if ((k & 0x7) == 0 && _mm256_testz_ps(cmp, cmp) != 0) {
+                        break;
+                    }
                 }
-                adder = _mm256_and_ps(adder, cmp);
-                counter = _mm256_add_ps(counter, adder);
-                if ((k & 0x7) == 0 && _mm256_testz_ps(cmp, cmp) != 0) {
-                    break;
+            }
+            else {
+                for (int k = 0; k < info.maxIter; k++) {
+                    __m256 aa = _mm256_mul_ps(a, a);
+                    __m256 bb = _mm256_mul_ps(b, b);
+                    __m256 abab = _mm256_mul_ps(a, b); abab = _mm256_add_ps(abab, abab);
+                    a = _mm256_add_ps(_mm256_sub_ps(aa, bb), xs);
+                    b = _mm256_add_ps(abab, ys);
+                    __m256 cmp = _mm256_cmp_ps(_mm256_add_ps(aa, bb), threshold, _CMP_LE_OQ);
+                    adder = _mm256_and_ps(adder, cmp);
+                    counter = _mm256_add_ps(counter, adder);
+                    if ((k & 0x7) == 0 && _mm256_testz_ps(cmp, cmp) != 0) {
+                        break;
+                    }
                 }
             }
 
+
             auto alignVec = [](float* data) -> float* {
                 void* aligned = data;
                 ::size_t length = 64;
@@ -133,21 +148,36 @@ void CpuGenerator<double, mnd::X86_AVX, parallel>::generate(const mnd::MandelInf
             __m256d a = xs;
             __m256d b = ys;
 
-            for (int k = 0; k < info.maxIter; k++) {
-                __m256d aa = _mm256_mul_pd(a, a);
-                __m256d bb = _mm256_mul_pd(b, b);
-                __m256d abab = _mm256_mul_pd(a, b); abab = _mm256_add_pd(abab, abab);
-                a = _mm256_add_pd(_mm256_sub_pd(aa, bb), xs);
-                b = _mm256_add_pd(abab, ys);
-                __m256d cmp = _mm256_cmp_pd(_mm256_add_pd(aa, bb), threshold, _CMP_LE_OQ);
-                if (info.smooth) {
+            if (info.smooth) {
+                for (int k = 0; k < info.maxIter; k++) {
+                    __m256d aa = _mm256_mul_pd(a, a);
+                    __m256d bb = _mm256_mul_pd(b, b);
+                    __m256d abab = _mm256_mul_pd(a, b); abab = _mm256_add_pd(abab, abab);
+                    a = _mm256_add_pd(_mm256_sub_pd(aa, bb), xs);
+                    b = _mm256_add_pd(abab, ys);
+                    __m256d cmp = _mm256_cmp_pd(_mm256_add_pd(aa, bb), threshold, _CMP_LE_OQ);
                     resultsa = _mm256_or_pd(_mm256_andnot_pd(cmp, resultsa), _mm256_and_pd(cmp, a));
                     resultsb = _mm256_or_pd(_mm256_andnot_pd(cmp, resultsb), _mm256_and_pd(cmp, b));
+                    adder = _mm256_and_pd(adder, cmp);
+                    counter = _mm256_add_pd(counter, adder);
+                    if ((k & 0x3) == 0 && _mm256_testz_si256(_mm256_castpd_si256(cmp), _mm256_castpd_si256(cmp)) != 0) {
+                        break;
+                    }
                 }
-                adder = _mm256_and_pd(adder, cmp);
-                counter = _mm256_add_pd(counter, adder);
-                if ((k & 0x3) == 0 && _mm256_testz_si256(_mm256_castpd_si256(cmp), _mm256_castpd_si256(cmp)) != 0) {
-                    break;
+            }
+            else {
+                for (int k = 0; k < info.maxIter; k++) {
+                    __m256d aa = _mm256_mul_pd(a, a);
+                    __m256d bb = _mm256_mul_pd(b, b);
+                    __m256d abab = _mm256_mul_pd(a, b); abab = _mm256_add_pd(abab, abab);
+                    a = _mm256_add_pd(_mm256_sub_pd(aa, bb), xs);
+                    b = _mm256_add_pd(abab, ys);
+                    __m256d cmp = _mm256_cmp_pd(_mm256_add_pd(aa, bb), threshold, _CMP_LE_OQ);
+                    adder = _mm256_and_pd(adder, cmp);
+                    counter = _mm256_add_pd(counter, adder);
+                    if ((k & 0x3) == 0 && _mm256_testz_si256(_mm256_castpd_si256(cmp), _mm256_castpd_si256(cmp)) != 0) {
+                        break;
+                    }
                 }
             }
 

+ 1 - 1
libmandel/src/Generators.cpp

@@ -96,7 +96,7 @@ namespace mnd
     {
         static const std::map<Precision, Real> precs {
             { Precision::FLOAT, getPrecision<float>() },
-            { Precision::DOUBLE_FLOAT, Real("1.0e-13") },
+            { Precision::DOUBLE_FLOAT, Real("4.0e-15") },
             { Precision::DOUBLE, getPrecision<double>() },
             { Precision::DOUBLE_DOUBLE, Real("1.0e-29") },
             { Precision::QUAD_DOUBLE, Real("1.0e-56") },

+ 4 - 2
libmandel/src/Mandel.cpp

@@ -186,6 +186,8 @@ std::unique_ptr<mnd::AdaptiveGenerator> MandelContext::createAdaptiveGenerator(v
         doubleGen = getCpuGenerator(GeneratorType::DOUBLE_SSE2);
     }
     if (cpuInfo.hasAvx() && cpuInfo.hasFma()) {
+        floatGen = getCpuGenerator(GeneratorType::FLOAT_AVX_FMA);
+        doubleGen = getCpuGenerator(GeneratorType::DOUBLE_AVX_FMA);
         doubleDoubleGen = getCpuGenerator(GeneratorType::DOUBLE_DOUBLE_AVX_FMA);
     }
 
@@ -259,7 +261,7 @@ std::vector<MandelDevice> MandelContext::createDevices(void)
             try {
                 md.generators.insert({ GeneratorType::FLOAT, std::make_unique<ClGeneratorFloat>(device) });
                 md.generators.insert({ GeneratorType::FIXED64, std::make_unique<ClGenerator64>(device) });
-                //md.generators.insert({ GeneratorType::DOUBLE_FLOAT, std::make_unique<ClGeneratorDoubleFloat>(device) });
+                md.generators.insert({ GeneratorType::DOUBLE_FLOAT, std::make_unique<ClGeneratorDoubleFloat>(device) });
             }
             catch (const std::string& err) {
                 printf("err: %s", err.c_str());
@@ -269,7 +271,7 @@ std::vector<MandelDevice> MandelContext::createDevices(void)
                 try {
                     md.generators.insert({ GeneratorType::DOUBLE, std::make_unique<ClGeneratorDouble>(device) });
                     md.generators.insert({ GeneratorType::DOUBLE_DOUBLE, std::make_unique<ClGeneratorDoubleDouble>(device) });
-                    md.generators.insert({ GeneratorType::QUAD_DOUBLE, std::make_unique<ClGeneratorQuadDouble>(device) });
+                    //md.generators.insert({ GeneratorType::QUAD_DOUBLE, std::make_unique<ClGeneratorQuadDouble>(device) });
                 }
                 catch (const std::string& err) {
                     printf("err: %s", err.c_str());

+ 11 - 0
libmandel/src/OpenClCode.cpp

@@ -1,11 +1,22 @@
 #include "OpenClCode.h"
 
+#include "opencl/float.h"
+#include "opencl/doublefloat.h"
+
 #include "opencl/fixed64.h"
 #include "opencl/fixed128.h"
 #include "opencl/fixed512.h"
 
 namespace mnd
 {
+    const char* getFloat_cl() {
+        return (char*) float_cl;
+    }
+
+    const char* getDoubleFloat_cl() {
+        return (char*) doublefloat_cl;
+    }
+
     const char* getFixed64_cl() {
         return (char*) fixed64_cl;
     }

+ 1 - 1
libmandel/src/Types.cpp

@@ -15,7 +15,7 @@ namespace mnd
         if (num == Real(0.0)) {
             return "0";
         }
-        int exponent = std::floor(static_cast<float>(mnd::log(num)) / ::logf(10.0));
+        int exponent = std::floor(static_cast<float>(mnd::log(num)) / ::logf(10.000001));
         float fac = static_cast<float>(num / mnd::pow(Real(10), Real(exponent)));
         std::stringstream ss;
         ss.precision(3);

+ 0 - 234
libmandel/src/doublefloat.h

@@ -1,234 +0,0 @@
-unsigned char doublefloat_cl[] = {
-  0x2f, 0x2f, 0x20, 0x63, 0x69, 0x74, 0x61, 0x74, 0x69, 0x6f, 0x6e, 0x3a,
-  0x20, 0x47, 0x75, 0x69, 0x6c, 0x6c, 0x61, 0x75, 0x6d, 0x65, 0x20, 0x64,
-  0x61, 0x20, 0x47, 0x72, 0x61, 0xc3, 0xa7, 0x63, 0x61, 0x2c, 0x20, 0x44,
-  0x61, 0x76, 0x69, 0x64, 0x20, 0x44, 0x65, 0x66, 0x6f, 0x75, 0x72, 0x2e,
-  0x20, 0x49, 0x6d, 0x70, 0x6c, 0x65, 0x6d, 0x65, 0x6e, 0x74, 0x61, 0x74,
-  0x69, 0x6f, 0x6e, 0x20, 0x6f, 0x66, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
-  0x2d, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6f, 0x70, 0x65, 0x72, 0x61,
-  0x74, 0x6f, 0x72, 0x73, 0x20, 0x6f, 0x6e, 0x20, 0x67, 0x72, 0x61, 0x70,
-  0x68, 0x69, 0x63, 0x73, 0x20, 0x68, 0x61, 0x72, 0x64, 0x77, 0x61, 0x72,
-  0x65, 0x2e, 0x0d, 0x0a, 0x2f, 0x2f, 0x20, 0x52, 0x65, 0x61, 0x6c, 0x20,
-  0x4e, 0x75, 0x6d, 0x62, 0x65, 0x72, 0x73, 0x20, 0x61, 0x6e, 0x64, 0x20,
-  0x43, 0x6f, 0x6d, 0x70, 0x75, 0x74, 0x65, 0x72, 0x73, 0x20, 0x37, 0x2c,
-  0x20, 0x4a, 0x75, 0x6c, 0x20, 0x32, 0x30, 0x30, 0x36, 0x2c, 0x20, 0x4e,
-  0x61, 0x6e, 0x63, 0x79, 0x2c, 0x20, 0x46, 0x72, 0x61, 0x6e, 0x63, 0x65,
-  0x2e, 0x20, 0x70, 0x70, 0x2e, 0x32, 0x33, 0x2d, 0x33, 0x32, 0x2e, 0x20,
-  0x66, 0x66, 0x68, 0x61, 0x6c, 0x2d, 0x30, 0x30, 0x30, 0x32, 0x31, 0x34,
-  0x34, 0x33, 0x0d, 0x0a, 0x2f, 0x2f, 0x20, 0x68, 0x74, 0x74, 0x70, 0x73,
-  0x3a, 0x2f, 0x2f, 0x68, 0x61, 0x6c, 0x2e, 0x61, 0x72, 0x63, 0x68, 0x69,
-  0x76, 0x65, 0x73, 0x2d, 0x6f, 0x75, 0x76, 0x65, 0x72, 0x74, 0x65, 0x73,
-  0x2e, 0x66, 0x72, 0x2f, 0x68, 0x61, 0x6c, 0x2d, 0x30, 0x30, 0x30, 0x32,
-  0x31, 0x34, 0x34, 0x33, 0x2f, 0x64, 0x6f, 0x63, 0x75, 0x6d, 0x65, 0x6e,
-  0x74, 0x0d, 0x0a, 0x0d, 0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20,
-  0x74, 0x77, 0x6f, 0x53, 0x75, 0x6d, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74,
-  0x20, 0x61, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62, 0x29,
-  0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x20, 0x73, 0x20, 0x3d, 0x20, 0x61, 0x20, 0x2b, 0x20, 0x62, 0x3b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
-  0x76, 0x20, 0x3d, 0x20, 0x73, 0x20, 0x2d, 0x20, 0x61, 0x3b, 0x0d, 0x0a,
-  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x20,
-  0x3d, 0x20, 0x28, 0x61, 0x20, 0x2d, 0x20, 0x28, 0x73, 0x20, 0x2d, 0x20,
-  0x76, 0x29, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x62, 0x20, 0x2d, 0x20, 0x76,
-  0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75,
-  0x72, 0x6e, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28,
-  0x73, 0x2c, 0x20, 0x72, 0x29, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a, 0x0d,
-  0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x73, 0x70, 0x6c, 0x69,
-  0x74, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x29, 0x20, 0x7b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
-  0x63, 0x20, 0x3d, 0x20, 0x28, 0x34, 0x30, 0x39, 0x36, 0x20, 0x2b, 0x20,
-  0x31, 0x29, 0x20, 0x2a, 0x20, 0x61, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20,
-  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x62, 0x69, 0x67, 0x20,
-  0x3d, 0x20, 0x63, 0x20, 0x2d, 0x20, 0x61, 0x3b, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x68, 0x69, 0x20,
-  0x3d, 0x20, 0x63, 0x20, 0x2d, 0x20, 0x61, 0x62, 0x69, 0x67, 0x3b, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61,
-  0x6c, 0x6f, 0x20, 0x3d, 0x20, 0x61, 0x20, 0x2d, 0x20, 0x61, 0x68, 0x69,
-  0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72,
-  0x6e, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x61,
-  0x68, 0x69, 0x2c, 0x20, 0x61, 0x6c, 0x6f, 0x29, 0x3b, 0x0d, 0x0a, 0x7d,
-  0x0d, 0x0a, 0x0d, 0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x74,
-  0x77, 0x6f, 0x50, 0x72, 0x6f, 0x64, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74,
-  0x20, 0x61, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62, 0x29,
-  0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x20, 0x78, 0x20, 0x3d, 0x20, 0x61, 0x20, 0x2a, 0x20, 0x62, 0x3b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
-  0x20, 0x61, 0x65, 0x78, 0x20, 0x3d, 0x20, 0x73, 0x70, 0x6c, 0x69, 0x74,
-  0x28, 0x61, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c,
-  0x6f, 0x61, 0x74, 0x32, 0x20, 0x62, 0x65, 0x78, 0x20, 0x3d, 0x20, 0x73,
-  0x70, 0x6c, 0x69, 0x74, 0x28, 0x62, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x65, 0x72, 0x72, 0x78,
-  0x20, 0x3d, 0x20, 0x78, 0x20, 0x2d, 0x20, 0x28, 0x61, 0x65, 0x78, 0x2e,
-  0x73, 0x30, 0x20, 0x2a, 0x20, 0x62, 0x65, 0x78, 0x2e, 0x73, 0x30, 0x29,
-  0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
-  0x20, 0x65, 0x72, 0x72, 0x79, 0x20, 0x3d, 0x20, 0x65, 0x72, 0x72, 0x78,
-  0x20, 0x2d, 0x20, 0x28, 0x61, 0x65, 0x78, 0x2e, 0x73, 0x31, 0x20, 0x2a,
-  0x20, 0x62, 0x65, 0x78, 0x2e, 0x73, 0x30, 0x29, 0x3b, 0x0d, 0x0a, 0x20,
-  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x65, 0x72, 0x72,
-  0x7a, 0x20, 0x3d, 0x20, 0x65, 0x72, 0x72, 0x79, 0x20, 0x2d, 0x20, 0x28,
-  0x61, 0x65, 0x78, 0x2e, 0x73, 0x30, 0x20, 0x2a, 0x20, 0x62, 0x65, 0x78,
-  0x2e, 0x73, 0x31, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66,
-  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x79, 0x20, 0x3d, 0x20, 0x28, 0x61, 0x65,
-  0x78, 0x2e, 0x73, 0x31, 0x20, 0x2a, 0x20, 0x62, 0x65, 0x78, 0x2e, 0x73,
-  0x31, 0x29, 0x20, 0x2d, 0x20, 0x65, 0x72, 0x72, 0x7a, 0x3b, 0x0d, 0x0a,
-  0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x28,
-  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x78, 0x2c, 0x20, 0x79,
-  0x29, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a, 0x0d, 0x0a, 0x66, 0x6c, 0x6f,
-  0x61, 0x74, 0x32, 0x20, 0x61, 0x64, 0x64, 0x28, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x32, 0x20, 0x61, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
-  0x20, 0x62, 0x29, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66,
-  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x20, 0x3d, 0x20, 0x61, 0x2e, 0x73,
-  0x30, 0x20, 0x2b, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x3b, 0x0d, 0x0a, 0x20,
-  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x73, 0x3b, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28, 0x66, 0x61, 0x62,
-  0x73, 0x28, 0x61, 0x2e, 0x73, 0x30, 0x29, 0x20, 0x3e, 0x3d, 0x20, 0x66,
-  0x61, 0x62, 0x73, 0x28, 0x62, 0x2e, 0x73, 0x30, 0x29, 0x29, 0x20, 0x7b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x73, 0x20,
-  0x3d, 0x20, 0x28, 0x28, 0x28, 0x61, 0x2e, 0x73, 0x30, 0x20, 0x2d, 0x20,
-  0x72, 0x29, 0x20, 0x2b, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x29, 0x20, 0x2b,
-  0x20, 0x62, 0x2e, 0x73, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x61, 0x2e, 0x73,
-  0x31, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0d, 0x0a, 0x20,
-  0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65, 0x20, 0x7b, 0x0d, 0x0a, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x73, 0x20, 0x3d, 0x20, 0x28,
-  0x28, 0x28, 0x62, 0x2e, 0x73, 0x30, 0x20, 0x2d, 0x20, 0x72, 0x29, 0x20,
-  0x2b, 0x20, 0x61, 0x2e, 0x73, 0x30, 0x29, 0x20, 0x2b, 0x20, 0x61, 0x2e,
-  0x73, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x62, 0x2e, 0x73, 0x31, 0x3b, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20,
-  0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x74, 0x77, 0x6f, 0x53, 0x75,
-  0x6d, 0x28, 0x72, 0x2c, 0x20, 0x73, 0x29, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d,
-  0x0a, 0x0d, 0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x6d, 0x75,
-  0x6c, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x61, 0x2c, 0x20,
-  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x62, 0x29, 0x20, 0x7b, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20,
-  0x74, 0x20, 0x3d, 0x20, 0x74, 0x77, 0x6f, 0x50, 0x72, 0x6f, 0x64, 0x28,
-  0x61, 0x2e, 0x73, 0x30, 0x2c, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x29, 0x3b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
-  0x74, 0x33, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x61, 0x2e, 0x73, 0x30, 0x20,
-  0x2a, 0x20, 0x62, 0x2e, 0x73, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x61,
-  0x2e, 0x73, 0x31, 0x20, 0x2a, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x29, 0x29,
-  0x20, 0x2b, 0x20, 0x74, 0x2e, 0x73, 0x31, 0x3b, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x74, 0x77, 0x6f,
-  0x53, 0x75, 0x6d, 0x28, 0x74, 0x2e, 0x73, 0x30, 0x2c, 0x20, 0x74, 0x2e,
-  0x73, 0x31, 0x29, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a, 0x0d, 0x0a, 0x66,
-  0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x6d, 0x75, 0x6c, 0x46, 0x6c, 0x6f,
-  0x61, 0x74, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x61, 0x2c,
-  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62, 0x29, 0x20, 0x7b, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20,
-  0x74, 0x20, 0x3d, 0x20, 0x74, 0x77, 0x6f, 0x50, 0x72, 0x6f, 0x64, 0x28,
-  0x61, 0x2e, 0x73, 0x30, 0x2c, 0x20, 0x62, 0x29, 0x3b, 0x0d, 0x0a, 0x20,
-  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x74, 0x33, 0x20,
-  0x3d, 0x20, 0x28, 0x61, 0x2e, 0x73, 0x31, 0x20, 0x2a, 0x20, 0x62, 0x29,
-  0x20, 0x2b, 0x20, 0x74, 0x2e, 0x73, 0x31, 0x3b, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x74, 0x77, 0x6f,
-  0x53, 0x75, 0x6d, 0x28, 0x74, 0x2e, 0x73, 0x30, 0x2c, 0x20, 0x74, 0x2e,
-  0x73, 0x31, 0x29, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a, 0x0d, 0x0a, 0x5f,
-  0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64,
-  0x20, 0x69, 0x74, 0x65, 0x72, 0x61, 0x74, 0x65, 0x28, 0x5f, 0x5f, 0x67,
-  0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a,
-  0x20, 0x41, 0x2c, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e,
-  0x74, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x2c, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x20, 0x78, 0x31, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
-  0x78, 0x32, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x79, 0x31,
-  0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x79, 0x32, 0x2c, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66,
-  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x70, 0x77, 0x31, 0x2c, 0x20, 0x66, 0x6c,
-  0x6f, 0x61, 0x74, 0x20, 0x70, 0x77, 0x32, 0x2c, 0x20, 0x66, 0x6c, 0x6f,
-  0x61, 0x74, 0x20, 0x70, 0x68, 0x31, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x20, 0x70, 0x68, 0x32, 0x2c, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x6d,
-  0x61, 0x78, 0x2c, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x73, 0x6d, 0x6f, 0x6f,
-  0x74, 0x68, 0x29, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x69,
-  0x6e, 0x74, 0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x3d, 0x20, 0x67,
-  0x65, 0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64,
-  0x28, 0x30, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x69, 0x6e,
-  0x74, 0x20, 0x70, 0x78, 0x20, 0x3d, 0x20, 0x69, 0x6e, 0x64, 0x65, 0x78,
-  0x20, 0x25, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x3b, 0x0d, 0x0a, 0x20,
-  0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x70, 0x79, 0x20, 0x3d, 0x20,
-  0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x2f, 0x20, 0x77, 0x69, 0x64, 0x74,
-  0x68, 0x3b, 0x0d, 0x0a, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c,
-  0x6f, 0x61, 0x74, 0x32, 0x20, 0x78, 0x6c, 0x20, 0x3d, 0x20, 0x28, 0x66,
-  0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x78, 0x31, 0x2c, 0x20, 0x78,
-  0x32, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f,
-  0x61, 0x74, 0x32, 0x20, 0x79, 0x74, 0x20, 0x3d, 0x20, 0x28, 0x66, 0x6c,
-  0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x79, 0x31, 0x2c, 0x20, 0x79, 0x32,
-  0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x32, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c,
-  0x65, 0x58, 0x20, 0x3d, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
-  0x29, 0x28, 0x70, 0x77, 0x31, 0x2c, 0x20, 0x70, 0x77, 0x32, 0x29, 0x3b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
-  0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x59,
-  0x20, 0x3d, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28,
-  0x70, 0x68, 0x31, 0x2c, 0x20, 0x70, 0x68, 0x32, 0x29, 0x3b, 0x0d, 0x0a,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
-  0x20, 0x61, 0x20, 0x3d, 0x20, 0x61, 0x64, 0x64, 0x28, 0x6d, 0x75, 0x6c,
-  0x46, 0x6c, 0x6f, 0x61, 0x74, 0x28, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53,
-  0x63, 0x61, 0x6c, 0x65, 0x58, 0x2c, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x29, 0x20, 0x70, 0x78, 0x29, 0x2c, 0x20, 0x78, 0x6c, 0x29, 0x3b,
-  0x20, 0x2f, 0x2f, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61,
-  0x6c, 0x65, 0x58, 0x20, 0x2a, 0x20, 0x70, 0x78, 0x20, 0x2b, 0x20, 0x78,
-  0x6c, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
-  0x32, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x61, 0x64, 0x64, 0x28, 0x6d, 0x75,
-  0x6c, 0x46, 0x6c, 0x6f, 0x61, 0x74, 0x28, 0x70, 0x69, 0x78, 0x65, 0x6c,
-  0x53, 0x63, 0x61, 0x6c, 0x65, 0x59, 0x2c, 0x20, 0x28, 0x66, 0x6c, 0x6f,
-  0x61, 0x74, 0x29, 0x20, 0x70, 0x79, 0x29, 0x2c, 0x20, 0x79, 0x74, 0x29,
-  0x3b, 0x20, 0x2f, 0x2f, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63,
-  0x61, 0x6c, 0x65, 0x59, 0x20, 0x2a, 0x20, 0x70, 0x79, 0x20, 0x2b, 0x20,
-  0x79, 0x74, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61,
-  0x74, 0x32, 0x20, 0x63, 0x61, 0x20, 0x3d, 0x20, 0x61, 0x3b, 0x0d, 0x0a,
-  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x63,
-  0x62, 0x20, 0x3d, 0x20, 0x62, 0x3b, 0x0d, 0x0a, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x6e, 0x20, 0x3d, 0x20, 0x30, 0x3b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x77, 0x68, 0x69, 0x6c, 0x65, 0x20,
-  0x28, 0x6e, 0x20, 0x3c, 0x20, 0x6d, 0x61, 0x78, 0x20, 0x2d, 0x20, 0x31,
-  0x29, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x61, 0x61, 0x20, 0x3d,
-  0x20, 0x6d, 0x75, 0x6c, 0x28, 0x61, 0x2c, 0x20, 0x61, 0x29, 0x3b, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f,
-  0x61, 0x74, 0x32, 0x20, 0x62, 0x62, 0x20, 0x3d, 0x20, 0x6d, 0x75, 0x6c,
-  0x28, 0x62, 0x2c, 0x20, 0x62, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20,
-  0x61, 0x62, 0x20, 0x3d, 0x20, 0x6d, 0x75, 0x6c, 0x28, 0x61, 0x2c, 0x20,
-  0x62, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x69, 0x66, 0x20, 0x28, 0x61, 0x61, 0x2e, 0x73, 0x30, 0x20, 0x2b,
-  0x20, 0x62, 0x62, 0x2e, 0x73, 0x30, 0x20, 0x3e, 0x20, 0x31, 0x36, 0x29,
-  0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20,
-  0x6d, 0x69, 0x6e, 0x75, 0x73, 0x62, 0x62, 0x20, 0x3d, 0x20, 0x28, 0x66,
-  0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x2d, 0x62, 0x62, 0x2e, 0x73,
-  0x30, 0x2c, 0x20, 0x2d, 0x62, 0x62, 0x2e, 0x73, 0x31, 0x29, 0x3b, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x61, 0x20, 0x3d,
-  0x20, 0x61, 0x64, 0x64, 0x28, 0x61, 0x64, 0x64, 0x28, 0x61, 0x61, 0x2c,
-  0x20, 0x6d, 0x69, 0x6e, 0x75, 0x73, 0x62, 0x62, 0x29, 0x2c, 0x20, 0x63,
-  0x61, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x62, 0x20, 0x3d, 0x20, 0x61, 0x64, 0x64, 0x28, 0x61, 0x64, 0x64,
-  0x28, 0x61, 0x62, 0x2c, 0x20, 0x61, 0x62, 0x29, 0x2c, 0x20, 0x63, 0x62,
-  0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x6e, 0x2b, 0x2b, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0d,
-  0x0a, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x2f, 0x2f, 0x20, 0x4e, 0x20,
-  0x2b, 0x20, 0x31, 0x20, 0x2d, 0x20, 0x6c, 0x6f, 0x67, 0x20, 0x28, 0x6c,
-  0x6f, 0x67, 0x20, 0x20, 0x7c, 0x5a, 0x28, 0x4e, 0x29, 0x7c, 0x29, 0x20,
-  0x2f, 0x20, 0x6c, 0x6f, 0x67, 0x20, 0x32, 0x0d, 0x0a, 0x20, 0x20, 0x20,
-  0x20, 0x69, 0x66, 0x20, 0x28, 0x6e, 0x20, 0x3e, 0x3d, 0x20, 0x6d, 0x61,
-  0x78, 0x20, 0x2d, 0x20, 0x31, 0x29, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x41, 0x5b, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x5d,
-  0x20, 0x3d, 0x20, 0x6d, 0x61, 0x78, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20,
-  0x20, 0x65, 0x6c, 0x73, 0x65, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28, 0x73, 0x6d, 0x6f,
-  0x6f, 0x74, 0x68, 0x20, 0x21, 0x3d, 0x20, 0x30, 0x29, 0x0d, 0x0a, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41,
-  0x5b, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x28,
-  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29, 0x20, 0x6e, 0x29, 0x20, 0x2b, 0x20,
-  0x31, 0x20, 0x2d, 0x20, 0x6c, 0x6f, 0x67, 0x28, 0x6c, 0x6f, 0x67, 0x28,
-  0x61, 0x2e, 0x73, 0x30, 0x20, 0x2a, 0x20, 0x61, 0x2e, 0x73, 0x30, 0x20,
-  0x2b, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x20, 0x2a, 0x20, 0x62, 0x2e, 0x73,
-  0x30, 0x29, 0x20, 0x2f, 0x20, 0x32, 0x29, 0x20, 0x2f, 0x20, 0x6c, 0x6f,
-  0x67, 0x28, 0x32, 0x2e, 0x30, 0x66, 0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65, 0x0d, 0x0a,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x41, 0x5b, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x5d, 0x20, 0x3d, 0x20, 0x28,
-  0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29, 0x6e, 0x29, 0x3b, 0x0d, 0x0a,
-  0x20, 0x20, 0x20, 0x20, 0x7d, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a
-};
-unsigned int doublefloat_cl_len = 2770;

+ 3 - 3
libmandel/src/doublefloat.cl → libmandel/src/opencl/doublefloat.cl

@@ -42,17 +42,17 @@ float2 add(float2 a, float2 b) {
 
 float2 mul(float2 a, float2 b) {
     float2 t = twoProd(a.s0, b.s0);
-    float t3 = ((a.s0 * b.s1) + (a.s1 * b.s0)) + t.s1;
+    t.s1 += ((a.s0 * b.s1) + (a.s1 * b.s0));
     return twoSum(t.s0, t.s1);
 }
 
 float2 mulFloat(float2 a, float b) {
     float2 t = twoProd(a.s0, b);
-    float t3 = (a.s1 * b) + t.s1;
+    t.s1 += (a.s1 * b);
     return twoSum(t.s0, t.s1);
 }
 
-__kernel void iterate(__global __write_only float* A, const int width,
+__kernel void iterate(__global float* A, const int width,
                       float x1, float x2, float y1, float y2,
                       float pw1, float pw2, float ph1, float ph2, int max, int smooth) {
     int index = get_global_id(0);

+ 225 - 0
libmandel/src/opencl/doublefloat.h

@@ -0,0 +1,225 @@
+unsigned char doublefloat_cl[] = {
+  0x2f, 0x2f, 0x20, 0x63, 0x69, 0x74, 0x61, 0x74, 0x69, 0x6f, 0x6e, 0x3a,
+  0x20, 0x47, 0x75, 0x69, 0x6c, 0x6c, 0x61, 0x75, 0x6d, 0x65, 0x20, 0x64,
+  0x61, 0x20, 0x47, 0x72, 0x61, 0xc3, 0xa7, 0x63, 0x61, 0x2c, 0x20, 0x44,
+  0x61, 0x76, 0x69, 0x64, 0x20, 0x44, 0x65, 0x66, 0x6f, 0x75, 0x72, 0x2e,
+  0x20, 0x49, 0x6d, 0x70, 0x6c, 0x65, 0x6d, 0x65, 0x6e, 0x74, 0x61, 0x74,
+  0x69, 0x6f, 0x6e, 0x20, 0x6f, 0x66, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x2d, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6f, 0x70, 0x65, 0x72, 0x61,
+  0x74, 0x6f, 0x72, 0x73, 0x20, 0x6f, 0x6e, 0x20, 0x67, 0x72, 0x61, 0x70,
+  0x68, 0x69, 0x63, 0x73, 0x20, 0x68, 0x61, 0x72, 0x64, 0x77, 0x61, 0x72,
+  0x65, 0x2e, 0x0a, 0x2f, 0x2f, 0x20, 0x52, 0x65, 0x61, 0x6c, 0x20, 0x4e,
+  0x75, 0x6d, 0x62, 0x65, 0x72, 0x73, 0x20, 0x61, 0x6e, 0x64, 0x20, 0x43,
+  0x6f, 0x6d, 0x70, 0x75, 0x74, 0x65, 0x72, 0x73, 0x20, 0x37, 0x2c, 0x20,
+  0x4a, 0x75, 0x6c, 0x20, 0x32, 0x30, 0x30, 0x36, 0x2c, 0x20, 0x4e, 0x61,
+  0x6e, 0x63, 0x79, 0x2c, 0x20, 0x46, 0x72, 0x61, 0x6e, 0x63, 0x65, 0x2e,
+  0x20, 0x70, 0x70, 0x2e, 0x32, 0x33, 0x2d, 0x33, 0x32, 0x2e, 0x20, 0x66,
+  0x66, 0x68, 0x61, 0x6c, 0x2d, 0x30, 0x30, 0x30, 0x32, 0x31, 0x34, 0x34,
+  0x33, 0x0a, 0x2f, 0x2f, 0x20, 0x68, 0x74, 0x74, 0x70, 0x73, 0x3a, 0x2f,
+  0x2f, 0x68, 0x61, 0x6c, 0x2e, 0x61, 0x72, 0x63, 0x68, 0x69, 0x76, 0x65,
+  0x73, 0x2d, 0x6f, 0x75, 0x76, 0x65, 0x72, 0x74, 0x65, 0x73, 0x2e, 0x66,
+  0x72, 0x2f, 0x68, 0x61, 0x6c, 0x2d, 0x30, 0x30, 0x30, 0x32, 0x31, 0x34,
+  0x34, 0x33, 0x2f, 0x64, 0x6f, 0x63, 0x75, 0x6d, 0x65, 0x6e, 0x74, 0x0a,
+  0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x74, 0x77, 0x6f, 0x53,
+  0x75, 0x6d, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x2c, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62, 0x29, 0x20, 0x7b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x73, 0x20, 0x3d,
+  0x20, 0x61, 0x20, 0x2b, 0x20, 0x62, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x76, 0x20, 0x3d, 0x20, 0x73, 0x20,
+  0x2d, 0x20, 0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x72, 0x20, 0x3d, 0x20, 0x28, 0x61, 0x20, 0x2d, 0x20,
+  0x28, 0x73, 0x20, 0x2d, 0x20, 0x76, 0x29, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x62, 0x20, 0x2d, 0x20, 0x76, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x32, 0x29, 0x28, 0x73, 0x2c, 0x20, 0x72, 0x29, 0x3b, 0x0a, 0x7d,
+  0x0a, 0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x73, 0x70, 0x6c,
+  0x69, 0x74, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x29, 0x20,
+  0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
+  0x63, 0x20, 0x3d, 0x20, 0x28, 0x34, 0x30, 0x39, 0x36, 0x20, 0x2b, 0x20,
+  0x31, 0x29, 0x20, 0x2a, 0x20, 0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x62, 0x69, 0x67, 0x20, 0x3d,
+  0x20, 0x63, 0x20, 0x2d, 0x20, 0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x68, 0x69, 0x20, 0x3d, 0x20,
+  0x63, 0x20, 0x2d, 0x20, 0x61, 0x62, 0x69, 0x67, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x6c, 0x6f, 0x20,
+  0x3d, 0x20, 0x61, 0x20, 0x2d, 0x20, 0x61, 0x68, 0x69, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x28, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x61, 0x68, 0x69, 0x2c, 0x20,
+  0x61, 0x6c, 0x6f, 0x29, 0x3b, 0x0a, 0x7d, 0x0a, 0x0a, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x32, 0x20, 0x74, 0x77, 0x6f, 0x50, 0x72, 0x6f, 0x64, 0x28,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x2c, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x62, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x78, 0x20, 0x3d, 0x20, 0x61, 0x20,
+  0x2a, 0x20, 0x62, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x32, 0x20, 0x61, 0x65, 0x78, 0x20, 0x3d, 0x20, 0x73, 0x70,
+  0x6c, 0x69, 0x74, 0x28, 0x61, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x62, 0x65, 0x78, 0x20, 0x3d,
+  0x20, 0x73, 0x70, 0x6c, 0x69, 0x74, 0x28, 0x62, 0x29, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x65, 0x72, 0x72,
+  0x78, 0x20, 0x3d, 0x20, 0x78, 0x20, 0x2d, 0x20, 0x28, 0x61, 0x65, 0x78,
+  0x2e, 0x73, 0x30, 0x20, 0x2a, 0x20, 0x62, 0x65, 0x78, 0x2e, 0x73, 0x30,
+  0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x65, 0x72, 0x72, 0x79, 0x20, 0x3d, 0x20, 0x65, 0x72, 0x72, 0x78,
+  0x20, 0x2d, 0x20, 0x28, 0x61, 0x65, 0x78, 0x2e, 0x73, 0x31, 0x20, 0x2a,
+  0x20, 0x62, 0x65, 0x78, 0x2e, 0x73, 0x30, 0x29, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x65, 0x72, 0x72, 0x7a,
+  0x20, 0x3d, 0x20, 0x65, 0x72, 0x72, 0x79, 0x20, 0x2d, 0x20, 0x28, 0x61,
+  0x65, 0x78, 0x2e, 0x73, 0x30, 0x20, 0x2a, 0x20, 0x62, 0x65, 0x78, 0x2e,
+  0x73, 0x31, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x79, 0x20, 0x3d, 0x20, 0x28, 0x61, 0x65, 0x78, 0x2e,
+  0x73, 0x31, 0x20, 0x2a, 0x20, 0x62, 0x65, 0x78, 0x2e, 0x73, 0x31, 0x29,
+  0x20, 0x2d, 0x20, 0x65, 0x72, 0x72, 0x7a, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x28, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x32, 0x29, 0x28, 0x78, 0x2c, 0x20, 0x79, 0x29, 0x3b, 0x0a,
+  0x7d, 0x0a, 0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x61, 0x64,
+  0x64, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x61, 0x2c, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x62, 0x29, 0x20, 0x7b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x20,
+  0x3d, 0x20, 0x61, 0x2e, 0x73, 0x30, 0x20, 0x2b, 0x20, 0x62, 0x2e, 0x73,
+  0x30, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28,
+  0x66, 0x61, 0x62, 0x73, 0x28, 0x61, 0x2e, 0x73, 0x30, 0x29, 0x20, 0x3e,
+  0x3d, 0x20, 0x66, 0x61, 0x62, 0x73, 0x28, 0x62, 0x2e, 0x73, 0x30, 0x29,
+  0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x73, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x28, 0x61, 0x2e, 0x73, 0x30, 0x20,
+  0x2d, 0x20, 0x72, 0x29, 0x20, 0x2b, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x29,
+  0x20, 0x2b, 0x20, 0x62, 0x2e, 0x73, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x61,
+  0x2e, 0x73, 0x31, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65, 0x20, 0x7b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x73, 0x20, 0x3d, 0x20, 0x28, 0x28,
+  0x28, 0x62, 0x2e, 0x73, 0x30, 0x20, 0x2d, 0x20, 0x72, 0x29, 0x20, 0x2b,
+  0x20, 0x61, 0x2e, 0x73, 0x30, 0x29, 0x20, 0x2b, 0x20, 0x61, 0x2e, 0x73,
+  0x31, 0x29, 0x20, 0x2b, 0x20, 0x62, 0x2e, 0x73, 0x31, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74,
+  0x75, 0x72, 0x6e, 0x20, 0x74, 0x77, 0x6f, 0x53, 0x75, 0x6d, 0x28, 0x72,
+  0x2c, 0x20, 0x73, 0x29, 0x3b, 0x0a, 0x7d, 0x0a, 0x0a, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x32, 0x20, 0x6d, 0x75, 0x6c, 0x28, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x32, 0x20, 0x61, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
+  0x20, 0x62, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x32, 0x20, 0x74, 0x20, 0x3d, 0x20, 0x74, 0x77, 0x6f,
+  0x50, 0x72, 0x6f, 0x64, 0x28, 0x61, 0x2e, 0x73, 0x30, 0x2c, 0x20, 0x62,
+  0x2e, 0x73, 0x30, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x74, 0x2e,
+  0x73, 0x31, 0x20, 0x2b, 0x3d, 0x20, 0x28, 0x28, 0x61, 0x2e, 0x73, 0x30,
+  0x20, 0x2a, 0x20, 0x62, 0x2e, 0x73, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x61, 0x2e, 0x73, 0x31, 0x20, 0x2a, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x29,
+  0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72,
+  0x6e, 0x20, 0x74, 0x77, 0x6f, 0x53, 0x75, 0x6d, 0x28, 0x74, 0x2e, 0x73,
+  0x30, 0x2c, 0x20, 0x74, 0x2e, 0x73, 0x31, 0x29, 0x3b, 0x0a, 0x7d, 0x0a,
+  0x0a, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x6d, 0x75, 0x6c, 0x46,
+  0x6c, 0x6f, 0x61, 0x74, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20,
+  0x61, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62, 0x29, 0x20,
+  0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
+  0x20, 0x74, 0x20, 0x3d, 0x20, 0x74, 0x77, 0x6f, 0x50, 0x72, 0x6f, 0x64,
+  0x28, 0x61, 0x2e, 0x73, 0x30, 0x2c, 0x20, 0x62, 0x29, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x74, 0x2e, 0x73, 0x31, 0x20, 0x2b, 0x3d, 0x20, 0x28,
+  0x61, 0x2e, 0x73, 0x31, 0x20, 0x2a, 0x20, 0x62, 0x29, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x74, 0x77,
+  0x6f, 0x53, 0x75, 0x6d, 0x28, 0x74, 0x2e, 0x73, 0x30, 0x2c, 0x20, 0x74,
+  0x2e, 0x73, 0x31, 0x29, 0x3b, 0x0a, 0x7d, 0x0a, 0x0a, 0x5f, 0x5f, 0x6b,
+  0x65, 0x72, 0x6e, 0x65, 0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x69,
+  0x74, 0x65, 0x72, 0x61, 0x74, 0x65, 0x28, 0x5f, 0x5f, 0x67, 0x6c, 0x6f,
+  0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x41,
+  0x2c, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20,
+  0x77, 0x69, 0x64, 0x74, 0x68, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x78,
+  0x31, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x78, 0x32, 0x2c,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x79, 0x31, 0x2c, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x79, 0x32, 0x2c, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x70, 0x77, 0x31, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
+  0x70, 0x77, 0x32, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x70,
+  0x68, 0x31, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x70, 0x68,
+  0x32, 0x2c, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x6d, 0x61, 0x78, 0x2c, 0x20,
+  0x69, 0x6e, 0x74, 0x20, 0x73, 0x6d, 0x6f, 0x6f, 0x74, 0x68, 0x29, 0x20,
+  0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x6e,
+  0x64, 0x65, 0x78, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f, 0x67, 0x6c,
+  0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x70, 0x78, 0x20, 0x3d,
+  0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x25, 0x20, 0x77, 0x69, 0x64,
+  0x74, 0x68, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20,
+  0x70, 0x79, 0x20, 0x3d, 0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x2f,
+  0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x78, 0x6c, 0x20, 0x3d,
+  0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x78, 0x31,
+  0x2c, 0x20, 0x78, 0x32, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x79, 0x74, 0x20, 0x3d, 0x20, 0x28,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x79, 0x31, 0x2c, 0x20,
+  0x79, 0x32, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x32, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61,
+  0x6c, 0x65, 0x58, 0x20, 0x3d, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x32, 0x29, 0x28, 0x70, 0x77, 0x31, 0x2c, 0x20, 0x70, 0x77, 0x32, 0x29,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
+  0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x59,
+  0x20, 0x3d, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x29, 0x28,
+  0x70, 0x68, 0x31, 0x2c, 0x20, 0x70, 0x68, 0x32, 0x29, 0x3b, 0x0a, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x61,
+  0x20, 0x3d, 0x20, 0x61, 0x64, 0x64, 0x28, 0x6d, 0x75, 0x6c, 0x46, 0x6c,
+  0x6f, 0x61, 0x74, 0x28, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61,
+  0x6c, 0x65, 0x58, 0x2c, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29,
+  0x20, 0x70, 0x78, 0x29, 0x2c, 0x20, 0x78, 0x6c, 0x29, 0x3b, 0x20, 0x2f,
+  0x2f, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65,
+  0x58, 0x20, 0x2a, 0x20, 0x70, 0x78, 0x20, 0x2b, 0x20, 0x78, 0x6c, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x62,
+  0x20, 0x3d, 0x20, 0x61, 0x64, 0x64, 0x28, 0x6d, 0x75, 0x6c, 0x46, 0x6c,
+  0x6f, 0x61, 0x74, 0x28, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61,
+  0x6c, 0x65, 0x59, 0x2c, 0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29,
+  0x20, 0x70, 0x79, 0x29, 0x2c, 0x20, 0x79, 0x74, 0x29, 0x3b, 0x20, 0x2f,
+  0x2f, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65,
+  0x59, 0x20, 0x2a, 0x20, 0x70, 0x79, 0x20, 0x2b, 0x20, 0x79, 0x74, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x63,
+  0x61, 0x20, 0x3d, 0x20, 0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x63, 0x62, 0x20, 0x3d, 0x20, 0x62,
+  0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x6e,
+  0x20, 0x3d, 0x20, 0x30, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x77, 0x68,
+  0x69, 0x6c, 0x65, 0x20, 0x28, 0x6e, 0x20, 0x3c, 0x20, 0x6d, 0x61, 0x78,
+  0x20, 0x2d, 0x20, 0x31, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x61,
+  0x61, 0x20, 0x3d, 0x20, 0x6d, 0x75, 0x6c, 0x28, 0x61, 0x2c, 0x20, 0x61,
+  0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x62, 0x62, 0x20, 0x3d, 0x20, 0x6d,
+  0x75, 0x6c, 0x28, 0x62, 0x2c, 0x20, 0x62, 0x29, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32,
+  0x20, 0x61, 0x62, 0x20, 0x3d, 0x20, 0x6d, 0x75, 0x6c, 0x28, 0x61, 0x2c,
+  0x20, 0x62, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x69, 0x66, 0x20, 0x28, 0x61, 0x61, 0x2e, 0x73, 0x30, 0x20, 0x2b,
+  0x20, 0x62, 0x62, 0x2e, 0x73, 0x30, 0x20, 0x3e, 0x20, 0x31, 0x36, 0x29,
+  0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x32, 0x20, 0x6d,
+  0x69, 0x6e, 0x75, 0x73, 0x62, 0x62, 0x20, 0x3d, 0x20, 0x28, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x32, 0x29, 0x28, 0x2d, 0x62, 0x62, 0x2e, 0x73, 0x30,
+  0x2c, 0x20, 0x2d, 0x62, 0x62, 0x2e, 0x73, 0x31, 0x29, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x61, 0x20, 0x3d, 0x20, 0x61,
+  0x64, 0x64, 0x28, 0x61, 0x64, 0x64, 0x28, 0x61, 0x61, 0x2c, 0x20, 0x6d,
+  0x69, 0x6e, 0x75, 0x73, 0x62, 0x62, 0x29, 0x2c, 0x20, 0x63, 0x61, 0x29,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x62, 0x20,
+  0x3d, 0x20, 0x61, 0x64, 0x64, 0x28, 0x61, 0x64, 0x64, 0x28, 0x61, 0x62,
+  0x2c, 0x20, 0x61, 0x62, 0x29, 0x2c, 0x20, 0x63, 0x62, 0x29, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x6e, 0x2b, 0x2b, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x2f, 0x2f, 0x20, 0x4e, 0x20, 0x2b, 0x20, 0x31, 0x20, 0x2d, 0x20, 0x6c,
+  0x6f, 0x67, 0x20, 0x28, 0x6c, 0x6f, 0x67, 0x20, 0x20, 0x7c, 0x5a, 0x28,
+  0x4e, 0x29, 0x7c, 0x29, 0x20, 0x2f, 0x20, 0x6c, 0x6f, 0x67, 0x20, 0x32,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28, 0x6e, 0x20, 0x3e,
+  0x3d, 0x20, 0x6d, 0x61, 0x78, 0x20, 0x2d, 0x20, 0x31, 0x29, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41, 0x5b, 0x69, 0x6e, 0x64,
+  0x65, 0x78, 0x5d, 0x20, 0x3d, 0x20, 0x6d, 0x61, 0x78, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65, 0x20, 0x7b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28, 0x73, 0x6d,
+  0x6f, 0x6f, 0x74, 0x68, 0x20, 0x21, 0x3d, 0x20, 0x30, 0x29, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41,
+  0x5b, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x28,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29, 0x20, 0x6e, 0x29, 0x20, 0x2b, 0x20,
+  0x31, 0x20, 0x2d, 0x20, 0x6c, 0x6f, 0x67, 0x28, 0x6c, 0x6f, 0x67, 0x28,
+  0x61, 0x2e, 0x73, 0x30, 0x20, 0x2a, 0x20, 0x61, 0x2e, 0x73, 0x30, 0x20,
+  0x2b, 0x20, 0x62, 0x2e, 0x73, 0x30, 0x20, 0x2a, 0x20, 0x62, 0x2e, 0x73,
+  0x30, 0x29, 0x20, 0x2f, 0x20, 0x32, 0x29, 0x20, 0x2f, 0x20, 0x6c, 0x6f,
+  0x67, 0x28, 0x32, 0x2e, 0x30, 0x66, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41, 0x5b,
+  0x69, 0x6e, 0x64, 0x65, 0x78, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x29, 0x6e, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x7d, 0x0a, 0x7d, 0x0a
+};
+unsigned int doublefloat_cl_len = 2657;

+ 67 - 0
libmandel/src/opencl/float.cl

@@ -0,0 +1,67 @@
+__kernel void iterate(__global float* A, const int width, float xl, float yt, float pixelScaleX, float pixelScaleY, int max, int smooth) {
+   int index = get_global_id(0);
+   int x = index % width;
+   int y = index / width;
+   float a = x * pixelScaleX + xl;
+   float b = y * pixelScaleY + yt;
+   float ca = a;
+   float cb = b;
+
+   int n = 0;
+   while (n < max - 1) {
+       float aa = a * a;
+       float bb = b * b;
+       float ab = a * b;
+       if (aa + bb > 16) break;
+       a = aa - bb + ca;
+       b = ab + ab + cb;
+       n++;
+   }
+   if (n >= max - 1) {
+       A[index] = max;
+   }
+   else {
+       if (smooth != 0)
+           A[index] = ((float)n) + 1 - log(log(a * a + b * b) / 2) / log(2.0f);
+       else
+           A[index] = ((float)n);
+   }
+}
+
+
+__kernel void iterate_vec4(__global float* A, const int width, float xl, float yt, float pixelScaleX, float pixelScaleY, int max, int smooth) {
+   int index = get_global_id(0) * 4;
+   int x = index % width;
+   int y = index / width;
+   float4 a = (float4) (x * pixelScaleX + xl, (x + 1) * pixelScaleX + xl, (x + 2) * pixelScaleX + xl, (x + 3) * pixelScaleX + xl);
+   float4 b = (float4) (y * pixelScaleY + yt);
+   float4 ca = a;
+   float4 cb = b;
+   float4 resa = a;
+   float4 resb = b;
+   int4 count = (int4)(0);
+
+   int n = 0;
+   while (n < max) {
+       float4 ab = a * b;
+       float4 cmpVal = fma(a, a, b * b);
+       int4 cmp = isless(cmpVal, (float4)(16.0f));
+       if (!any(cmp)) break;
+       a = fma(a, a, -fma(b, b, -ca));
+       b = fma(2, ab, cb);
+       if (smooth) {
+           resa = as_float4(as_int4(a) & cmp | (as_int4(resa) & ~cmp));
+           resb = as_float4(as_int4(b) & cmp | (as_int4(resb) & ~cmp));
+       }
+       count += cmp & (int4)(1);
+       n++;
+   }
+   for (int i = 0; i < 4 && i + x < width; i++) {
+       if (smooth != 0)
+           A[index + i] = ((float) count[i]) + 1 - log(log(fma(resa[i], resa[i], resb[i] * resb[i])) / 2) / log(2.0f);
+      else
+          A[index + i] = ((float) count[i]);
+   }
+}
+
+

+ 171 - 0
libmandel/src/opencl/float.h

@@ -0,0 +1,171 @@
+unsigned char float_cl[] = {
+  0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x20, 0x76, 0x6f, 0x69,
+  0x64, 0x20, 0x69, 0x74, 0x65, 0x72, 0x61, 0x74, 0x65, 0x28, 0x5f, 0x5f,
+  0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x2a, 0x20, 0x41, 0x2c, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69,
+  0x6e, 0x74, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x2c, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x78, 0x6c, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x20, 0x79, 0x74, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
+  0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x58, 0x2c,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c,
+  0x53, 0x63, 0x61, 0x6c, 0x65, 0x59, 0x2c, 0x20, 0x69, 0x6e, 0x74, 0x20,
+  0x6d, 0x61, 0x78, 0x2c, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x73, 0x6d, 0x6f,
+  0x6f, 0x74, 0x68, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e,
+  0x74, 0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x3d, 0x20, 0x67, 0x65,
+  0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28,
+  0x30, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x78,
+  0x20, 0x3d, 0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x25, 0x20, 0x77,
+  0x69, 0x64, 0x74, 0x68, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x79, 0x20, 0x3d, 0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x2f,
+  0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x20, 0x3d, 0x20, 0x78, 0x20, 0x2a,
+  0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x58,
+  0x20, 0x2b, 0x20, 0x78, 0x6c, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x79, 0x20, 0x2a, 0x20,
+  0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x59, 0x20,
+  0x2b, 0x20, 0x79, 0x74, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x63, 0x61, 0x20, 0x3d, 0x20, 0x61, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x63, 0x62, 0x20, 0x3d,
+  0x20, 0x62, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20,
+  0x6e, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x77, 0x68,
+  0x69, 0x6c, 0x65, 0x20, 0x28, 0x6e, 0x20, 0x3c, 0x20, 0x6d, 0x61, 0x78,
+  0x20, 0x2d, 0x20, 0x31, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x61, 0x20,
+  0x3d, 0x20, 0x61, 0x20, 0x2a, 0x20, 0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62, 0x62,
+  0x20, 0x3d, 0x20, 0x62, 0x20, 0x2a, 0x20, 0x62, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61,
+  0x62, 0x20, 0x3d, 0x20, 0x61, 0x20, 0x2a, 0x20, 0x62, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28, 0x61, 0x61,
+  0x20, 0x2b, 0x20, 0x62, 0x62, 0x20, 0x3e, 0x20, 0x31, 0x36, 0x29, 0x20,
+  0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x61, 0x20, 0x3d, 0x20, 0x61, 0x61, 0x20, 0x2d, 0x20, 0x62,
+  0x62, 0x20, 0x2b, 0x20, 0x63, 0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x61, 0x62, 0x20, 0x2b, 0x20,
+  0x61, 0x62, 0x20, 0x2b, 0x20, 0x63, 0x62, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x6e, 0x2b, 0x2b, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x7d, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28, 0x6e, 0x20, 0x3e,
+  0x3d, 0x20, 0x6d, 0x61, 0x78, 0x20, 0x2d, 0x20, 0x31, 0x29, 0x20, 0x7b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41, 0x5b, 0x69, 0x6e,
+  0x64, 0x65, 0x78, 0x5d, 0x20, 0x3d, 0x20, 0x6d, 0x61, 0x78, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65,
+  0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66,
+  0x20, 0x28, 0x73, 0x6d, 0x6f, 0x6f, 0x74, 0x68, 0x20, 0x21, 0x3d, 0x20,
+  0x30, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x41, 0x5b, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x5d, 0x20, 0x3d,
+  0x20, 0x28, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29, 0x6e, 0x29, 0x20,
+  0x2b, 0x20, 0x31, 0x20, 0x2d, 0x20, 0x6c, 0x6f, 0x67, 0x28, 0x6c, 0x6f,
+  0x67, 0x28, 0x61, 0x20, 0x2a, 0x20, 0x61, 0x20, 0x2b, 0x20, 0x62, 0x20,
+  0x2a, 0x20, 0x62, 0x29, 0x20, 0x2f, 0x20, 0x32, 0x29, 0x20, 0x2f, 0x20,
+  0x6c, 0x6f, 0x67, 0x28, 0x32, 0x2e, 0x30, 0x66, 0x29, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41, 0x5b,
+  0x69, 0x6e, 0x64, 0x65, 0x78, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x29, 0x6e, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x7d, 0x0a, 0x7d, 0x0a, 0x0a, 0x0a, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e,
+  0x65, 0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x69, 0x74, 0x65, 0x72,
+  0x61, 0x74, 0x65, 0x5f, 0x76, 0x65, 0x63, 0x34, 0x28, 0x5f, 0x5f, 0x67,
+  0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a,
+  0x20, 0x41, 0x2c, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e,
+  0x74, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x2c, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x78, 0x6c, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x79, 0x74, 0x2c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x70,
+  0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x58, 0x2c, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53,
+  0x63, 0x61, 0x6c, 0x65, 0x59, 0x2c, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x6d,
+  0x61, 0x78, 0x2c, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x73, 0x6d, 0x6f, 0x6f,
+  0x74, 0x68, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74,
+  0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30,
+  0x29, 0x20, 0x2a, 0x20, 0x34, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e,
+  0x74, 0x20, 0x78, 0x20, 0x3d, 0x20, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20,
+  0x25, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x69, 0x6e, 0x74, 0x20, 0x79, 0x20, 0x3d, 0x20, 0x69, 0x6e, 0x64, 0x65,
+  0x78, 0x20, 0x2f, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34, 0x20, 0x61, 0x20, 0x3d,
+  0x20, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34, 0x29, 0x20, 0x28, 0x78,
+  0x20, 0x2a, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c,
+  0x65, 0x58, 0x20, 0x2b, 0x20, 0x78, 0x6c, 0x2c, 0x20, 0x28, 0x78, 0x20,
+  0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c,
+  0x53, 0x63, 0x61, 0x6c, 0x65, 0x58, 0x20, 0x2b, 0x20, 0x78, 0x6c, 0x2c,
+  0x20, 0x28, 0x78, 0x20, 0x2b, 0x20, 0x32, 0x29, 0x20, 0x2a, 0x20, 0x70,
+  0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x58, 0x20, 0x2b,
+  0x20, 0x78, 0x6c, 0x2c, 0x20, 0x28, 0x78, 0x20, 0x2b, 0x20, 0x33, 0x29,
+  0x20, 0x2a, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c,
+  0x65, 0x58, 0x20, 0x2b, 0x20, 0x78, 0x6c, 0x29, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34, 0x20, 0x62, 0x20, 0x3d, 0x20,
+  0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34, 0x29, 0x20, 0x28, 0x79, 0x20,
+  0x2a, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65,
+  0x59, 0x20, 0x2b, 0x20, 0x79, 0x74, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34, 0x20, 0x63, 0x61, 0x20, 0x3d, 0x20,
+  0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34,
+  0x20, 0x63, 0x62, 0x20, 0x3d, 0x20, 0x62, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34, 0x20, 0x72, 0x65, 0x73, 0x61, 0x20,
+  0x3d, 0x20, 0x61, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x34, 0x20, 0x72, 0x65, 0x73, 0x62, 0x20, 0x3d, 0x20, 0x62, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x34, 0x20, 0x63, 0x6f, 0x75,
+  0x6e, 0x74, 0x20, 0x3d, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x34, 0x29, 0x28,
+  0x30, 0x29, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x20,
+  0x6e, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x77, 0x68,
+  0x69, 0x6c, 0x65, 0x20, 0x28, 0x6e, 0x20, 0x3c, 0x20, 0x6d, 0x61, 0x78,
+  0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x34, 0x20, 0x61, 0x62, 0x20, 0x3d, 0x20, 0x61,
+  0x20, 0x2a, 0x20, 0x62, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x34, 0x20, 0x63, 0x6d, 0x70, 0x56,
+  0x61, 0x6c, 0x20, 0x3d, 0x20, 0x66, 0x6d, 0x61, 0x28, 0x61, 0x2c, 0x20,
+  0x61, 0x2c, 0x20, 0x62, 0x20, 0x2a, 0x20, 0x62, 0x29, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x6e, 0x74, 0x34, 0x20, 0x63,
+  0x6d, 0x70, 0x20, 0x3d, 0x20, 0x69, 0x73, 0x6c, 0x65, 0x73, 0x73, 0x28,
+  0x63, 0x6d, 0x70, 0x56, 0x61, 0x6c, 0x2c, 0x20, 0x28, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x34, 0x29, 0x28, 0x31, 0x36, 0x2e, 0x30, 0x66, 0x29, 0x29,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20,
+  0x28, 0x21, 0x61, 0x6e, 0x79, 0x28, 0x63, 0x6d, 0x70, 0x29, 0x29, 0x20,
+  0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x61, 0x20, 0x3d, 0x20, 0x66, 0x6d, 0x61, 0x28, 0x61, 0x2c,
+  0x20, 0x61, 0x2c, 0x20, 0x2d, 0x66, 0x6d, 0x61, 0x28, 0x62, 0x2c, 0x20,
+  0x62, 0x2c, 0x20, 0x2d, 0x63, 0x61, 0x29, 0x29, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x66, 0x6d, 0x61,
+  0x28, 0x32, 0x2c, 0x20, 0x61, 0x62, 0x2c, 0x20, 0x63, 0x62, 0x29, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28,
+  0x73, 0x6d, 0x6f, 0x6f, 0x74, 0x68, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x73,
+  0x61, 0x20, 0x3d, 0x20, 0x61, 0x73, 0x5f, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x34, 0x28, 0x61, 0x73, 0x5f, 0x69, 0x6e, 0x74, 0x34, 0x28, 0x61, 0x29,
+  0x20, 0x26, 0x20, 0x63, 0x6d, 0x70, 0x20, 0x7c, 0x20, 0x28, 0x61, 0x73,
+  0x5f, 0x69, 0x6e, 0x74, 0x34, 0x28, 0x72, 0x65, 0x73, 0x61, 0x29, 0x20,
+  0x26, 0x20, 0x7e, 0x63, 0x6d, 0x70, 0x29, 0x29, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x73,
+  0x62, 0x20, 0x3d, 0x20, 0x61, 0x73, 0x5f, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x34, 0x28, 0x61, 0x73, 0x5f, 0x69, 0x6e, 0x74, 0x34, 0x28, 0x62, 0x29,
+  0x20, 0x26, 0x20, 0x63, 0x6d, 0x70, 0x20, 0x7c, 0x20, 0x28, 0x61, 0x73,
+  0x5f, 0x69, 0x6e, 0x74, 0x34, 0x28, 0x72, 0x65, 0x73, 0x62, 0x29, 0x20,
+  0x26, 0x20, 0x7e, 0x63, 0x6d, 0x70, 0x29, 0x29, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x75, 0x6e, 0x74, 0x20, 0x2b, 0x3d, 0x20, 0x63,
+  0x6d, 0x70, 0x20, 0x26, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x34, 0x29, 0x28,
+  0x31, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x6e,
+  0x2b, 0x2b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20,
+  0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d,
+  0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x34, 0x20, 0x26, 0x26,
+  0x20, 0x69, 0x20, 0x2b, 0x20, 0x78, 0x20, 0x3c, 0x20, 0x77, 0x69, 0x64,
+  0x74, 0x68, 0x3b, 0x20, 0x69, 0x2b, 0x2b, 0x29, 0x20, 0x7b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x69, 0x66, 0x20, 0x28, 0x73, 0x6d,
+  0x6f, 0x6f, 0x74, 0x68, 0x20, 0x21, 0x3d, 0x20, 0x30, 0x29, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41, 0x5b,
+  0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x2b, 0x20, 0x69, 0x5d, 0x20, 0x3d,
+  0x20, 0x28, 0x28, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29, 0x20, 0x63, 0x6f,
+  0x75, 0x6e, 0x74, 0x5b, 0x69, 0x5d, 0x29, 0x20, 0x2b, 0x20, 0x31, 0x20,
+  0x2d, 0x20, 0x6c, 0x6f, 0x67, 0x28, 0x6c, 0x6f, 0x67, 0x28, 0x66, 0x6d,
+  0x61, 0x28, 0x72, 0x65, 0x73, 0x61, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x72,
+  0x65, 0x73, 0x61, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x72, 0x65, 0x73, 0x62,
+  0x5b, 0x69, 0x5d, 0x20, 0x2a, 0x20, 0x72, 0x65, 0x73, 0x62, 0x5b, 0x69,
+  0x5d, 0x29, 0x29, 0x20, 0x2f, 0x20, 0x32, 0x29, 0x20, 0x2f, 0x20, 0x6c,
+  0x6f, 0x67, 0x28, 0x32, 0x2e, 0x30, 0x66, 0x29, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x65, 0x6c, 0x73, 0x65, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x41, 0x5b, 0x69, 0x6e, 0x64,
+  0x65, 0x78, 0x20, 0x2b, 0x20, 0x69, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x28,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29, 0x20, 0x63, 0x6f, 0x75, 0x6e, 0x74,
+  0x5b, 0x69, 0x5d, 0x29, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x7d,
+  0x0a, 0x0a, 0x0a
+};
+unsigned int float_cl_len = 2007;