opencl fixed64 works (not yet feingeschliffen)

Gradient.cpp

@@ -96,7 +96,7 @@ Gradient Gradient::readXml(const QString& xml)
 
 RGBColor Gradient::get(float x) const
 {
-    if (colors.empty() || ::isnan(x) || ::isinf(x))
+    if (colors.empty() || std::isnan(x) || std::isinf(x))
         return RGBColor();
     /*const auto [left, right, lerp] = getNeighbors(x);
     RGBColor lerped = lerpColors(left, right, lerp);

libmandel/include/Fixed.h

@@ -534,34 +534,35 @@ public:
 
 struct Fixed64
 {
-    bool sign;
-    uint64_t bits;
+    int64_t bits;
 
     Fixed64(const Fixed64&) = default;
     ~Fixed64() = default;
 
 
-    inline Fixed64(uint64_t bits, bool /* dummy */) :
+    inline Fixed64(int64_t bits, bool /* dummy */) :
         bits{ bits }
     {
     }
 
     inline Fixed64(double x)
     {
-        if (x < 0) {
-            sign = true;
-            x *= -1;
-        }
-        else {
-            sign = false;
-        }
-        int integerPart = int(x);
+        /*int integerPart = int(x);
         double fractionalPart = x - integerPart;
         bits = uint64_t(integerPart) << 32;
         bits |= uint64_t(fractionalPart * (1ULL << 32)) & 0xFFFFFFFF;
+        */
+        bits = x * (1LL << 32);
+    }
+
+    inline operator float(void) const {
+        return bits * (1.0f / (1ULL << 32));
+    }
+    inline operator double(void) const {
+        return bits * (1.0 / (1ULL << 32));
     }
 
-    inline Fixed64 operator + (const Fixed64& other) {
+    inline Fixed64 operator + (const Fixed64& other) const {
         return Fixed64{ bits + other.bits, true };
     }
     
@@ -570,9 +571,49 @@ struct Fixed64
         return *this;
     }
 
-    inline Fixed64 operator - (const Fixed64& other) {
+    inline Fixed64 operator - (const Fixed64& other) const {
         return Fixed64{ bits - other.bits, true };
     }
+    inline Fixed64& operator -= (const Fixed64& other) {
+        bits -= other.bits;
+        return *this;
+    }
+
+    static inline std::pair<int64_t, uint64_t> mul64(int64_t a, int64_t b) {
+        uint32_t aa[2] = { uint32_t(a >> 32), uint32_t(a & 0xFFFFFFFF) };
+        uint32_t bb[2] = { uint32_t(b >> 32), uint32_t(b & 0xFFFFFFFF) };
+
+        uint32_t res[4];
+        int64_t temp = uint64_t(aa[1]) * uint64_t(bb[1]);
+        res[3] = temp & 0xFFFFFFFF;
+        temp >>= 32;
+        temp += int64_t(int32_t(aa[0])) * int64_t(bb[1]) + int64_t(aa[1]) * int64_t(int32_t(bb[0]));
+        res[2] = temp & 0xFFFFFFFF;
+        temp >>= 32;
+        temp += int64_t(int32_t(aa[0])) * int64_t(int32_t(bb[0]));
+        res[1] = temp & 0xFFFFFFFF;
+        res[0] = temp >> 32;
+
+        return std::make_pair((int64_t(res[0]) << 32) | res[1], uint64_t((int64_t(res[2]) << 32) | res[3]));
+    }
+
+    static inline std::pair<uint64_t, uint64_t> mulu64(uint64_t a, uint64_t b) {
+        uint32_t aa[2] = { uint32_t(a >> 32), uint32_t(a & 0xFFFFFFFF) };
+        uint32_t bb[2] = { uint32_t(b >> 32), uint32_t(b & 0xFFFFFFFF) };
+
+        uint32_t res[4];
+        uint64_t temp = uint64_t(aa[1]) * bb[1];
+        res[3] = temp & 0xFFFFFFFF;
+        uint32_t carry = temp >> 32;
+        temp = uint64_t(aa[0]) * bb[1] + uint64_t(aa[1]) * bb[0] + carry;
+        res[2] = temp & 0xFFFFFFFF;
+        carry = temp >> 32;
+        temp = uint64_t(aa[0]) * bb[0] + carry;
+        res[1] = temp & 0xFFFFFFFF;
+        res[0] = temp >> 32;
+
+        return std::make_pair((uint64_t(res[0]) << 32) | res[1], (uint64_t(res[2]) << 32) | res[3] );
+    }
 
     inline Fixed64 operator * (const Fixed64& other) {
         /*int32_t upper = bits >> 32;
@@ -583,10 +624,44 @@ struct Fixed64
 
         int32_t newUp = upup & 0xFFFFFFFF + (loup >> 32);
         int32_t newLo = loup & 0xFFFFFFFF + (lolo >> 32);*/
-        double d = int32_t(bits >> 32) + double(uint32_t(bits)) / (1ULL << 32);
+        /*double d = int32_t(bits >> 32) + double(uint32_t(bits)) / (1ULL << 32);
         double od = int32_t(other.bits >> 32) + double(uint32_t(other.bits)) / (1ULL << 32);
-        return d * od * ((other.sign != sign) ? -1 : 1);
+        return d * od;*/
+
+        /*auto[hi, lo] = mul64(bits, other.bits);
+        return Fixed64{ int64_t((hi << 32) | (lo >> 32)), true };*/
+
+        uint32_t a[2] = { uint32_t(uint64_t(bits) >> 32), uint32_t(bits & 0xFFFFFFFF) };
+        uint32_t b[2] = { uint32_t(uint64_t(other.bits) >> 32), uint32_t(other.bits & 0xFFFFFFFF) };
+
+        uint64_t a1b1 = uint64_t(a[1]) * b[1];
+        int64_t a0b1 = int64_t(int32_t(a[0])) * uint64_t(b[1]);
+        int64_t a1b0 = uint64_t(a[1]) * int64_t(int32_t(b[1]));
+        int64_t a0b0 = int64_t(int32_t(a[1])) * int64_t(int32_t(b[1]));
+
+        int64_t res = a1b1 >> 32;
+        res += a0b1 + a1b0;
+        res += a0b0 << 32;
+        return Fixed64{ res, true };
 
+        /*
+        uint32_t aa[2] = { uint32_t(uint64_t(bits) >> 32), uint32_t(bits & 0xFFFFFFFF) };
+        uint32_t bb[2] = { uint32_t(uint64_t(other.bits) >> 32), uint32_t(other.bits & 0xFFFFFFFF) };
+
+        uint64_t ab0[2] = { 0, 0 };
+        ab[1] = int64_t(int32_t(aa[1]) * int32_t(ab[0]));
+        ab[0] = int64_t(int32_t(aa[0]) * int32_t(ab[0]));
+
+        uint64_t ab1[2] = { 0, 0 };
+        ab[1] = aa[1] * ab[1];
+        ab[0] = int64_t(int32_t(aa[1]) * int32_t(ab[0]));
+        */
+
+        /*
+        boost::multiprecision::int128_t a(this->bits);
+        boost::multiprecision::int128_t b(other.bits);
+        return Fixed64{ int64_t((a * b) >> 32), true };
+        */
         //return Fixed64{ (uint64_t(newUp) << 32) | newLo, true };
     }
 

libmandel/include/Types.h

@@ -137,6 +137,12 @@ namespace mnd
     {
         return float(Real(x));
     }
+
+    template<>
+    inline Fixed64 convert<Fixed64, Float512>(const Float512& x)
+    {
+        return Fixed64(double(Real(x)));
+    }
 #endif
 
     std::string toString(const Real& num);

libmandel/src/ClGenerators.cpp

@@ -621,6 +621,7 @@ ClGenerator64::ClGenerator64(cl::Device device) :
 }
 
 
+#include "CpuGenerators.h"
 void ClGenerator64::generate(const mnd::MandelInfo& info, float* data)
 {
     ::size_t bufferSize = info.bWidth * info.bHeight * sizeof(float);
@@ -630,11 +631,14 @@ void ClGenerator64::generate(const mnd::MandelInfo& info, float* data)
     float pixelScaleY = float(info.view.height / info.bHeight);
 
     using ull = unsigned long long;
-    ull x = ull(double(info.view.x) * 0x1000000000000ULL);
-    ull y = ull(double(info.view.y) * 0x1000000000000ULL);
-    ull w = ull(double(pixelScaleX) * 0x1000000000000ULL);
-    ull h = ull(double(pixelScaleY) * 0x1000000000000ULL);
+    ull x = ull(::round(double(info.view.x) * (1LL << 32)));
+    ull y = ull(::round(double(info.view.y) * (1LL << 32)));
+    ull w = ull(::round(double(pixelScaleX) * (1LL << 32)));
+    ull h = ull(::round(double(pixelScaleY) * (1LL << 32)));
+    //x = 0;
+    //y = 0;
 
+    
     Kernel iterate = Kernel(program, "iterate");
     iterate.setArg(0, buffer_A);
     iterate.setArg(1, int(info.bWidth));
@@ -647,6 +651,8 @@ void ClGenerator64::generate(const mnd::MandelInfo& info, float* data)
 
     queue.enqueueNDRangeKernel(iterate, 0, NDRange(info.bWidth * info.bHeight));
     queue.enqueueReadBuffer(buffer_A, CL_TRUE, 0, bufferSize, data);
+    //CpuGenerator<Fixed64> fx;
+    //fx.generate(info, data);
 }
 
 

libmandel/src/CpuGenerators.cpp

@@ -20,6 +20,8 @@ namespace mnd
     template class CpuGenerator<double, mnd::NONE, false>;
     template class CpuGenerator<double, mnd::NONE, true>;
 
+    template class CpuGenerator<Fixed64, mnd::NONE, false>;
+    template class CpuGenerator<Fixed64, mnd::NONE, true>;
     
     //template class CpuGenerator<Fixed128, mnd::NONE, false>;
     //template class CpuGenerator<Fixed128, mnd::NONE, false, true>;

libmandel/src/opencl/fixed64.cl

@@ -10,16 +10,16 @@ long mul(long a, long b) {
 __kernel void iterate(__global float* A, const int width,
                       ulong x, ulong y, ulong pw, ulong ph, int max, int smooth) {
     int index = get_global_id(0);
-    long px = (index % width) << 48;
-    long py = (index / width) << 48;
+    long px = (index % width) ;
+    long py = (index / width) ;
 
     long xl = x;
     long yt = y;
     long pixelScaleX = pw;
     long pixelScaleY = ph;
 
-    long a = (mul(pixelScaleX, px) + xl); // pixelScaleX * px + xl
-    long b = (mul(pixelScaleY, py) + yt); // pixelScaleY * py + yt
+    long a = xl + pixelScaleX * px; // pixelScaleX * px + xl
+    long b = yt + pixelScaleY * py; // pixelScaleY * py + yt
     long ca = a;
     long cb = b;
 
@@ -28,7 +28,7 @@ __kernel void iterate(__global float* A, const int width,
         long aa = mul(a, a);
         long bb = mul(b, b);
         long ab = mul(a, b);
-        if (aa + bb > (16LL << 48)) break;
+        if (aa + bb > (16LL << 32)) break;
         a = aa - bb + ca;
         b = ab + ab + cb;
         n++;
@@ -39,13 +39,11 @@ __kernel void iterate(__global float* A, const int width,
         A[index] = max;
     else {
         if (smooth != 0) {
-            float aapprox = ((float) a) * 3.5527137e-15f;
-            float bapprox = ((float) b) * 3.5527137e-15f;
+            float aapprox = ((float) a) * (1.0f / (1LL << 32)); // 3.5527137e-15f;
+            float bapprox = ((float) b) * (1.0f / (1LL << 32)); // 3.5527137e-15f;
             A[index] = ((float) n) + 1 - log(log(aapprox * aapprox + bapprox * bapprox) / 2) / log(2.0f);
         }
         else
             A[index] = ((float)n);
     }
-    //               A[index] = ((float)n) + 1 - (a * a + b * b - 16) / (256 - 16);
-    //           A[get_global_id(0)] = 5;
 }

libmandel/src/opencl/fixed64.h

@@ -1,127 +1,115 @@
 unsigned char fixed64_cl[] = {
-  0x0d, 0x0a, 0x0d, 0x0a, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x6d, 0x75, 0x6c,
-  0x28, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x61, 0x2c, 0x20, 0x6c, 0x6f, 0x6e,
-  0x67, 0x20, 0x62, 0x29, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20,
-  0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x75, 0x70, 0x70, 0x65, 0x72, 0x20, 0x3d,
-  0x20, 0x6d, 0x75, 0x6c, 0x5f, 0x68, 0x69, 0x28, 0x61, 0x2c, 0x20, 0x62,
-  0x29, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67,
-  0x20, 0x6c, 0x6f, 0x77, 0x65, 0x72, 0x20, 0x3d, 0x20, 0x61, 0x20, 0x2a,
-  0x20, 0x62, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74,
-  0x75, 0x72, 0x6e, 0x20, 0x28, 0x75, 0x70, 0x70, 0x65, 0x72, 0x20, 0x3c,
-  0x3c, 0x20, 0x33, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x6c, 0x6f,
-  0x77, 0x65, 0x72, 0x20, 0x3e, 0x3e, 0x20, 0x33, 0x32, 0x29, 0x20, 0x26,
-  0x20, 0x30, 0x78, 0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x29,
-  0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a, 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,
+  0x0a, 0x0a, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x6d, 0x75, 0x6c, 0x28, 0x6c,
+  0x6f, 0x6e, 0x67, 0x20, 0x61, 0x2c, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20,
+  0x62, 0x29, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e,
+  0x67, 0x20, 0x75, 0x70, 0x70, 0x65, 0x72, 0x20, 0x3d, 0x20, 0x6d, 0x75,
+  0x6c, 0x5f, 0x68, 0x69, 0x28, 0x61, 0x2c, 0x20, 0x62, 0x29, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x6c, 0x6f, 0x77,
+  0x65, 0x72, 0x20, 0x3d, 0x20, 0x61, 0x20, 0x2a, 0x20, 0x62, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x28,
+  0x75, 0x70, 0x70, 0x65, 0x72, 0x20, 0x3c, 0x3c, 0x20, 0x33, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x28, 0x6c, 0x6f, 0x77, 0x65, 0x72, 0x20, 0x3e,
+  0x3e, 0x20, 0x33, 0x32, 0x29, 0x20, 0x26, 0x20, 0x30, 0x78, 0x46, 0x46,
+  0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x29, 0x3b, 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, 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, 0x75, 0x6c, 0x6f, 0x6e, 0x67,
-  0x20, 0x78, 0x2c, 0x20, 0x75, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x79, 0x2c,
-  0x20, 0x75, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70, 0x77, 0x2c, 0x20, 0x75,
-  0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70, 0x68, 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,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x75, 0x6c, 0x6f,
+  0x6e, 0x67, 0x20, 0x78, 0x2c, 0x20, 0x75, 0x6c, 0x6f, 0x6e, 0x67, 0x20,
+  0x79, 0x2c, 0x20, 0x75, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70, 0x77, 0x2c,
+  0x20, 0x75, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70, 0x68, 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,
   0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70, 0x78, 0x20, 0x3d, 0x20, 0x28, 0x69,
   0x6e, 0x64, 0x65, 0x78, 0x20, 0x25, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68,
-  0x29, 0x20, 0x3c, 0x3c, 0x20, 0x34, 0x38, 0x3b, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70, 0x79, 0x20, 0x3d, 0x20,
-  0x28, 0x69, 0x6e, 0x64, 0x65, 0x78, 0x20, 0x2f, 0x20, 0x77, 0x69, 0x64,
-  0x74, 0x68, 0x29, 0x20, 0x3c, 0x3c, 0x20, 0x34, 0x38, 0x3b, 0x0d, 0x0a,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x78,
-  0x6c, 0x20, 0x3d, 0x20, 0x78, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20,
-  0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x79, 0x74, 0x20, 0x3d, 0x20, 0x79, 0x3b,
-  0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70,
-  0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x58, 0x20, 0x3d,
-  0x20, 0x70, 0x77, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f,
-  0x6e, 0x67, 0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c,
-  0x65, 0x59, 0x20, 0x3d, 0x20, 0x70, 0x68, 0x3b, 0x0d, 0x0a, 0x0d, 0x0a,
-  0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x61, 0x20, 0x3d,
-  0x20, 0x28, 0x6d, 0x75, 0x6c, 0x28, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53,
-  0x63, 0x61, 0x6c, 0x65, 0x58, 0x2c, 0x20, 0x70, 0x78, 0x29, 0x20, 0x2b,
-  0x20, 0x78, 0x6c, 0x29, 0x3b, 0x20, 0x2f, 0x2f, 0x20, 0x70, 0x69, 0x78,
+  0x29, 0x20, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67,
+  0x20, 0x70, 0x79, 0x20, 0x3d, 0x20, 0x28, 0x69, 0x6e, 0x64, 0x65, 0x78,
+  0x20, 0x2f, 0x20, 0x77, 0x69, 0x64, 0x74, 0x68, 0x29, 0x20, 0x3b, 0x0a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x78, 0x6c,
+  0x20, 0x3d, 0x20, 0x78, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f,
+  0x6e, 0x67, 0x20, 0x79, 0x74, 0x20, 0x3d, 0x20, 0x79, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70, 0x69, 0x78, 0x65,
+  0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x58, 0x20, 0x3d, 0x20, 0x70, 0x77,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x70,
+  0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x59, 0x20, 0x3d,
+  0x20, 0x70, 0x68, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f,
+  0x6e, 0x67, 0x20, 0x61, 0x20, 0x3d, 0x20, 0x78, 0x6c, 0x20, 0x2b, 0x20,
+  0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x58, 0x20,
+  0x2a, 0x20, 0x70, 0x78, 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,
-  0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x28, 0x6d, 0x75,
-  0x6c, 0x28, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65,
-  0x59, 0x2c, 0x20, 0x70, 0x79, 0x29, 0x20, 0x2b, 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, 0x6c, 0x6f, 0x6e, 0x67,
-  0x20, 0x63, 0x61, 0x20, 0x3d, 0x20, 0x61, 0x3b, 0x0d, 0x0a, 0x20, 0x20,
-  0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 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,
+  0x78, 0x20, 0x2b, 0x20, 0x78, 0x6c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c,
+  0x6f, 0x6e, 0x67, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x79, 0x74, 0x20, 0x2b,
+  0x20, 0x70, 0x69, 0x78, 0x65, 0x6c, 0x53, 0x63, 0x61, 0x6c, 0x65, 0x59,
+  0x20, 0x2a, 0x20, 0x70, 0x79, 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,
+  0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x63, 0x61, 0x20, 0x3d, 0x20, 0x61, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67, 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, 0x0d,
-  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e,
-  0x67, 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, 0x6c, 0x6f, 0x6e, 0x67, 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, 0x6c, 0x6f, 0x6e,
-  0x67, 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, 0x20, 0x2b, 0x20,
-  0x62, 0x62, 0x20, 0x3e, 0x20, 0x28, 0x31, 0x36, 0x4c, 0x4c, 0x20, 0x3c,
-  0x3c, 0x20, 0x34, 0x38, 0x29, 0x29, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b,
-  0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x61,
-  0x20, 0x3d, 0x20, 0x61, 0x61, 0x20, 0x2d, 0x20, 0x62, 0x62, 0x20, 0x2b,
-  0x20, 0x63, 0x61, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x61, 0x62, 0x20, 0x2b, 0x20, 0x61,
-  0x62, 0x20, 0x2b, 0x20, 0x63, 0x62, 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,
+  0x20, 0x6d, 0x61, 0x78, 0x20, 0x2d, 0x20, 0x31, 0x29, 0x20, 0x7b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x6c, 0x6f, 0x6e, 0x67,
+  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, 0x6c, 0x6f, 0x6e, 0x67, 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, 0x6c, 0x6f, 0x6e, 0x67, 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, 0x20, 0x2b, 0x20, 0x62, 0x62, 0x20, 0x3e,
+  0x20, 0x28, 0x31, 0x36, 0x4c, 0x4c, 0x20, 0x3c, 0x3c, 0x20, 0x33, 0x32,
+  0x29, 0x29, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 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, 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, 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,
-  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, 0x20, 0x7b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
-  0x61, 0x61, 0x70, 0x70, 0x72, 0x6f, 0x78, 0x20, 0x3d, 0x20, 0x28, 0x28,
-  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x29, 0x20, 0x61, 0x29, 0x20, 0x2a, 0x20,
-  0x33, 0x2e, 0x35, 0x35, 0x32, 0x37, 0x31, 0x33, 0x37, 0x65, 0x2d, 0x31,
-  0x35, 0x66, 0x3b, 0x0d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62,
-  0x61, 0x70, 0x70, 0x72, 0x6f, 0x78, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x66,
-  0x6c, 0x6f, 0x61, 0x74, 0x29, 0x20, 0x62, 0x29, 0x20, 0x2a, 0x20, 0x33,
-  0x2e, 0x35, 0x35, 0x32, 0x37, 0x31, 0x33, 0x37, 0x65, 0x2d, 0x31, 0x35,
-  0x66, 0x3b, 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, 0x61, 0x70, 0x70, 0x72, 0x6f, 0x78,
-  0x20, 0x2a, 0x20, 0x61, 0x61, 0x70, 0x70, 0x72, 0x6f, 0x78, 0x20, 0x2b,
-  0x20, 0x62, 0x61, 0x70, 0x70, 0x72, 0x6f, 0x78, 0x20, 0x2a, 0x20, 0x62,
-  0x61, 0x70, 0x70, 0x72, 0x6f, 0x78, 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, 0x7d,
-  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, 0x20,
-  0x20, 0x20, 0x20, 0x2f, 0x2f, 0x20, 0x20, 0x20, 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, 0x20, 0x2b, 0x20, 0x31, 0x20, 0x2d, 0x20,
-  0x28, 0x61, 0x20, 0x2a, 0x20, 0x61, 0x20, 0x2b, 0x20, 0x62, 0x20, 0x2a,
-  0x20, 0x62, 0x20, 0x2d, 0x20, 0x31, 0x36, 0x29, 0x20, 0x2f, 0x20, 0x28,
-  0x32, 0x35, 0x36, 0x20, 0x2d, 0x20, 0x31, 0x36, 0x29, 0x3b, 0x0d, 0x0a,
-  0x20, 0x20, 0x20, 0x20, 0x2f, 0x2f, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
-  0x20, 0x20, 0x20, 0x20, 0x20, 0x41, 0x5b, 0x67, 0x65, 0x74, 0x5f, 0x67,
-  0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29, 0x5d,
-  0x20, 0x3d, 0x20, 0x35, 0x3b, 0x0d, 0x0a, 0x7d, 0x0d, 0x0a
+  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,
+  0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x61, 0x61, 0x70,
+  0x70, 0x72, 0x6f, 0x78, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x29, 0x20, 0x61, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x31, 0x2e,
+  0x30, 0x66, 0x20, 0x2f, 0x20, 0x28, 0x31, 0x4c, 0x4c, 0x20, 0x3c, 0x3c,
+  0x20, 0x33, 0x32, 0x29, 0x29, 0x3b, 0x20, 0x2f, 0x2f, 0x20, 0x33, 0x2e,
+  0x35, 0x35, 0x32, 0x37, 0x31, 0x33, 0x37, 0x65, 0x2d, 0x31, 0x35, 0x66,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x62, 0x61, 0x70, 0x70,
+  0x72, 0x6f, 0x78, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x29, 0x20, 0x62, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x31, 0x2e, 0x30,
+  0x66, 0x20, 0x2f, 0x20, 0x28, 0x31, 0x4c, 0x4c, 0x20, 0x3c, 0x3c, 0x20,
+  0x33, 0x32, 0x29, 0x29, 0x3b, 0x20, 0x2f, 0x2f, 0x20, 0x33, 0x2e, 0x35,
+  0x35, 0x32, 0x37, 0x31, 0x33, 0x37, 0x65, 0x2d, 0x31, 0x35, 0x66, 0x3b,
+  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, 0x61, 0x70, 0x70, 0x72, 0x6f, 0x78, 0x20, 0x2a, 0x20,
+  0x61, 0x61, 0x70, 0x70, 0x72, 0x6f, 0x78, 0x20, 0x2b, 0x20, 0x62, 0x61,
+  0x70, 0x70, 0x72, 0x6f, 0x78, 0x20, 0x2a, 0x20, 0x62, 0x61, 0x70, 0x70,
+  0x72, 0x6f, 0x78, 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, 0x7d, 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 fixed64_cl_len = 1486;
+unsigned int fixed64_cl_len = 1337;