]> foleosoft.com Git - QAnsel.git/commitdiff
Thu Mar 21 01:32:40 AM EDT 2024
authormiha-q <>
Thu, 21 Mar 2024 05:32:40 +0000 (01:32 -0400)
committermiha-q <>
Thu, 21 Mar 2024 05:32:40 +0000 (01:32 -0400)
50 files changed:
Makefile
bin/QAnsel
build.sh
examples/belltest_c.txt [deleted file]
examples/belltest_q.txt [deleted file]
examples/chsh_inequality_c.txt [new file with mode: 0644]
examples/chsh_inequality_q.txt [new file with mode: 0644]
examples/slow.txt
obj/bytecode.o [new file with mode: 0644]
obj/complex.o [new file with mode: 0644]
obj/context.o [new file with mode: 0644]
obj/display.o [new file with mode: 0644]
obj/hardware.o [new file with mode: 0644]
obj/kernel_cpu.o [new file with mode: 0644]
obj/main.o [new file with mode: 0644]
obj/openqasm.o [new file with mode: 0644]
src/.kernel.tmp.1 [new file with mode: 0644]
src/.kernel.tmp.2 [new file with mode: 0644]
src/bytecode.c
src/bytecode.h [new file with mode: 0644]
src/bytecode.o [new file with mode: 0644]
src/complex.c
src/complex.h [new file with mode: 0644]
src/complex.o [new file with mode: 0644]
src/context.c
src/context.h [new file with mode: 0644]
src/context.o [new file with mode: 0644]
src/display.c
src/display.h [new file with mode: 0644]
src/display.o [new file with mode: 0644]
src/finaldone [new file with mode: 0755]
src/gates.c [deleted file]
src/hardware.c
src/hardware.h [new file with mode: 0644]
src/hardware.o [new file with mode: 0644]
src/interface.html [deleted file]
src/kernel.cl
src/kernel.h [new file with mode: 0644]
src/kernel_cpu.c [new file with mode: 0644]
src/kernel_cpu.cl [new file with mode: 0644]
src/kernel_cpu.o [new file with mode: 0644]
src/kernel_gpu.c [new file with mode: 0644]
src/kernel_gpu.cl [new file with mode: 0644]
src/main.c
src/main.h [new file with mode: 0644]
src/main.o [new file with mode: 0644]
src/openqasm.c
src/openqasm.h [new file with mode: 0644]
src/openqasm.o [new file with mode: 0644]
src/qansel.h

index 29c84d50fde75ee5817ad70cddc00a7951953d99..3683099713c1cf0668f4188f20af46b2453d5057 100644 (file)
--- a/Makefile
+++ b/Makefile
@@ -4,6 +4,9 @@ all:
 simple:
        sh build.sh simple
 
+cl:
+       sh build.sh cl
+
 commit:
        git add src/ examples/ Makefile build.sh
        git commit -m "`date`"
index 347663844e56a383667a53a3863b25cce2464905..f6ac158f043c105a25dd0835276ddb033820d61d 100755 (executable)
Binary files a/bin/QAnsel and b/bin/QAnsel differ
index f2c500dbb4fa15440ffd722a4a6214af6b13836a..47d8c03f094fcb985e52cc06d349710c775e0f16 100644 (file)
--- a/build.sh
+++ b/build.sh
@@ -1,7 +1,5 @@
 #!/bin/sh
 
-echo "$1"
-
 echo "Verifying build commands exist..."
 for i in cat grep sed xxd gcc rm printf uname
 do
@@ -13,27 +11,49 @@ do
 done
 echo "Verified."
 echo "Building..."
-cat src/kernel.cl | grep -vi '{gpu_only}' | sed -e 's/__global //' -e 's/__kernel //' > src/kernel_cpu.cl
+cat src/kernel.cl | grep -vi '{gpu_only}' | sed -e 's/__global //' -e 's/__kernel //' > src/kernel_cpu.c
 cat src/kernel.cl | grep -vi '{cpu_only}' > src/.kernel.tmp.1
 tmp="$(cat src/.kernel.tmp.1)"
 printf "%s\0" "$tmp" > src/.kernel.tmp.2
-xxd -i src/.kernel.tmp.2 | sed -e 's/src__kernel_tmp_2/kernel_gpu/' > src/kernel_gpu.cl
+xxd -i src/.kernel.tmp.2 | sed -e 's/src__kernel_tmp_2/kernel_gpu/' -e 's/unsigned/static unsigned/' > src/kernel_gpu.c
+if [ "$1" = "cl" ]
+then
+       exit
+fi
+
+#check if x86_64
 x86=""
 if [ "$(uname -m)" = "x86_64" ]
 then
        x86=" -D__x86_64__"
 fi
+
 if [ "$1" = "simple" ]
 then
-       bcmd="gcc src/main.c -g -o bin/QAnsel -lm$x86"
-else
-       bcmd="gcc src/main.c -g -o bin/QAnsel -lm$x86 -D__PTHREAD__ -D__SDL2__ -D__OPENCL__ -I/usr/include/SDL2 -D_REENTRANT -lSDL2 -lOpenCL -pthread"
-fi
-echo "$bcmd"
-if $($bcmd)
-then
-       echo "Built."
+       cflags="-I/usr/include/SDL2"
 else
-       echo "Failed."
+       cflags="-I/usr/include/SDL2 -D__PTHREAD__ -D__SDL2__ -D__OPENCL__ -I/usr/include/SDL2 -D_REENTRANT"
 fi
-rm -f src/.kernel* src/kernel_*
\ No newline at end of file
+
+cd src/
+
+rm -r ../obj 2>/dev/null
+mkdir ../obj 2>/dev/null
+
+fcmd="gcc -o ../bin/QAnsel"
+for i in complex context display hardware bytecode openqasm main
+do
+       bcmd="gcc $i.c -c -o ../obj/$i.o $cflags"
+       echo "$bcmd"
+       $($bcmd)
+       fcmd="$fcmd $i.o" 
+done
+bcmd="gcc kernel_cpu.c -c -o ../obj/kernel_cpu.o $cflags" && echo "$bcmd" && $($bcmd)
+echo $bcmd
+$($bcmd)
+
+
+cd ../obj
+fcmd="$fcmd kernel_cpu.o -lm -D_REENTRANT -lSDL2 -lOpenCL -pthread" 
+echo "$fcmd"
+$($fcmd)
diff --git a/examples/belltest_c.txt b/examples/belltest_c.txt
deleted file mode 100644 (file)
index 8218c24..0000000
+++ /dev/null
@@ -1,56 +0,0 @@
-qreg q[3];
-creg c[4];
-
-//Classic Strategy
-
-//Host generates random bits
-//  for the two players.
-h q[0];
-h q[1];
-measure q[0] -> c[0];
-measure q[1] -> c[1];
-barrier q;
-
-//Player X strategy
-x q[0];
-
-//Player Y strategy
-
-//Transfer to host
-measure q[0] -> c[2];
-measure q[1] -> c[3];
-barrier q;
-
-//a xor b
-cx q[0], q[1];
-measure q[1] -> c[3];
-barrier q;
-
-//Load x and y
-reset q[0];
-
-if(c[0]==1) x q[0];
-reset q[1];
-if(c[1]==1) x q[1];
-barrier q;
-
-//x and y
-ccx q[0], q[1], q[2];
-measure q[2] -> c[2];
-barrier q;
-
-//Load (a xor b) and (x and y)
-reset q[0];
-if(c[2]==1) x q[0];
-reset q[1];
-if(c[3]==1) x q[1];
-barrier q;
-
-//(a xor b) = (x and y)
-cx q[0], q[1];
-x q[1];
-barrier q;
-
-//Sample results
-measure q[1] -> c[0];
-sample c[0];
\ No newline at end of file
diff --git a/examples/belltest_q.txt b/examples/belltest_q.txt
deleted file mode 100644 (file)
index 8dab919..0000000
+++ /dev/null
@@ -1,67 +0,0 @@
-qreg q[3];
-creg c[4];
-
-//Host generates random bits
-//  for the two players.
-h q[0];
-h q[1];
-measure q[0] -> c[0];
-measure q[1] -> c[1];
-reset q[0];
-reset q[1];
-barrier q;
-
-//Two players are also provided
-//  an entangled qubit.
-h q[0];
-cx q[0], q[1];
-barrier q;
-
-//Player X strategy
-//if(c[0]==0) do nothing
-if(c[0]==1) ry(pi/2) q[0];
-barrier q;
-
-//Player Y strategy
-if(c[1]==0) ry(pi/4) q[1];
-if(c[1]==1) ry(-pi/4) q[1];
-barrier q;
-
-//Transfer to host
-measure q[0] -> c[2];
-measure q[1] -> c[3];
-barrier q;
-
-//a xor b
-cx q[0], q[1];
-measure q[1] -> c[3];
-barrier q;
-
-//Load x and y
-reset q[0];
-if(c[0]==1) x q[0];
-reset q[1];
-if(c[1]==1) x q[1];
-barrier q;
-
-//x and y
-ccx q[0], q[1], q[2];
-measure q[2] -> c[2];
-barrier q;
-
-//Load (a xor b) and (x and y)
-reset q[0];
-if(c[2]==1) x q[0];
-reset q[1];
-if(c[3]==1) x q[1];
-barrier q;
-
-//(a xor b) = (x and y)
-cx q[0], q[1];
-x q[1];
-barrier q;
-
-//Store final results
-measure q[1] -> c[0];
-
-sample c[0];
\ No newline at end of file
diff --git a/examples/chsh_inequality_c.txt b/examples/chsh_inequality_c.txt
new file mode 100644 (file)
index 0000000..8218c24
--- /dev/null
@@ -0,0 +1,56 @@
+qreg q[3];
+creg c[4];
+
+//Classic Strategy
+
+//Host generates random bits
+//  for the two players.
+h q[0];
+h q[1];
+measure q[0] -> c[0];
+measure q[1] -> c[1];
+barrier q;
+
+//Player X strategy
+x q[0];
+
+//Player Y strategy
+
+//Transfer to host
+measure q[0] -> c[2];
+measure q[1] -> c[3];
+barrier q;
+
+//a xor b
+cx q[0], q[1];
+measure q[1] -> c[3];
+barrier q;
+
+//Load x and y
+reset q[0];
+
+if(c[0]==1) x q[0];
+reset q[1];
+if(c[1]==1) x q[1];
+barrier q;
+
+//x and y
+ccx q[0], q[1], q[2];
+measure q[2] -> c[2];
+barrier q;
+
+//Load (a xor b) and (x and y)
+reset q[0];
+if(c[2]==1) x q[0];
+reset q[1];
+if(c[3]==1) x q[1];
+barrier q;
+
+//(a xor b) = (x and y)
+cx q[0], q[1];
+x q[1];
+barrier q;
+
+//Sample results
+measure q[1] -> c[0];
+sample c[0];
\ No newline at end of file
diff --git a/examples/chsh_inequality_q.txt b/examples/chsh_inequality_q.txt
new file mode 100644 (file)
index 0000000..815e7c2
--- /dev/null
@@ -0,0 +1,57 @@
+//CHSH Inequality Bell Test
+qreg q[3];
+creg c[4];
+
+//Host generates random bits
+//  for the two players.
+h q[0];
+h q[1];
+measure q[0] -> c[0];
+measure q[1] -> c[1];
+reset q[0];
+reset q[1];
+
+//Two players are also provided
+//  an entangled qubit.
+h q[0];
+cx q[0], q[1];
+
+//Player X strategy
+//if(c[0]==0) do nothing
+if(c[0]==1) ry(pi/2) q[0];
+
+//Player Y strategy
+if(c[1]==0) ry(pi/4) q[1];
+if(c[1]==1) ry(-pi/4) q[1];
+
+//Transfer to host
+measure q[0] -> c[2];
+measure q[1] -> c[3];
+
+//a xor b
+cx q[0], q[1];
+measure q[1] -> c[3];
+
+//Load x and y
+reset q[0];
+if(c[0]==1) x q[0];
+reset q[1];
+if(c[1]==1) x q[1];
+
+//x and y
+ccx q[0], q[1], q[2];
+measure q[2] -> c[2];
+
+//Load (a xor b) and (x and y)
+reset q[0];
+if(c[2]==1) x q[0];
+reset q[1];
+if(c[3]==1) x q[1];
+
+//(a xor b) = (x and y)
+cx q[0], q[1];
+x q[1];
+
+//Store final results
+measure q[1] -> c[0];
+sample c[0];
index 65eb3a99c4709412e8caada2c3bb49b1528fa5b8..63975126d4b9daf4935b7c5dd032cb9dd56681cc 100644 (file)
@@ -1,20 +1,3 @@
-//This normally runs really slow and thus is a good
-//     illustration of how to use various optimization
-//     settings.
-//     -o 0            No optimization
-//     -o 1            Threads 
-//     -o 2            GPU
-//     -o 3            Threads+GPU
-//     -o 4            Queue
-//     -o 5            Queue+Threads
-//     -o 6            Queue+GPU
-//     -o 7            Queue+Threads+GPU
-//
-//The command queue will group together all
-//     consecutive single-qubit instructions. This
-//     program will run almost as fast as if there is
-//     just a single instruction.
-
 qreg q[12];
 h q[0];
 x q[1];
@@ -28,4 +11,4 @@ x q[8];
 x q[9];
 x q[10];
 x q[11];
-born q;
\ No newline at end of file
+born q;
diff --git a/obj/bytecode.o b/obj/bytecode.o
new file mode 100644 (file)
index 0000000..4379a3f
Binary files /dev/null and b/obj/bytecode.o differ
diff --git a/obj/complex.o b/obj/complex.o
new file mode 100644 (file)
index 0000000..41776dd
Binary files /dev/null and b/obj/complex.o differ
diff --git a/obj/context.o b/obj/context.o
new file mode 100644 (file)
index 0000000..83f557e
Binary files /dev/null and b/obj/context.o differ
diff --git a/obj/display.o b/obj/display.o
new file mode 100644 (file)
index 0000000..8e7c04b
Binary files /dev/null and b/obj/display.o differ
diff --git a/obj/hardware.o b/obj/hardware.o
new file mode 100644 (file)
index 0000000..9057ffc
Binary files /dev/null and b/obj/hardware.o differ
diff --git a/obj/kernel_cpu.o b/obj/kernel_cpu.o
new file mode 100644 (file)
index 0000000..baa933a
Binary files /dev/null and b/obj/kernel_cpu.o differ
diff --git a/obj/main.o b/obj/main.o
new file mode 100644 (file)
index 0000000..9397f00
Binary files /dev/null and b/obj/main.o differ
diff --git a/obj/openqasm.o b/obj/openqasm.o
new file mode 100644 (file)
index 0000000..46135d2
Binary files /dev/null and b/obj/openqasm.o differ
diff --git a/src/.kernel.tmp.1 b/src/.kernel.tmp.1
new file mode 100644 (file)
index 0000000..52960f9
--- /dev/null
@@ -0,0 +1,145 @@
+
+__kernel void kernel_dot
+(
+    __global float* ptrR,
+    __global float* ptrA,
+    __global float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB
+)
+{
+    const int rowsR = rowsA;
+    const int colsR = colsB;
+    const int rowR = get_global_id(0); //{gpu_only}
+    const int colR = get_global_id(1); //{gpu_only}
+    
+    float rR = 0;
+    float iR = 0;
+
+    for (int i = 0; i < colsA; i++)
+    {
+        const float rA = ptrA[(size_t)rowR * ((size_t)colsA * (size_t)2) + ((size_t)i * (size_t)2)];
+        const float iA = ptrA[(size_t)rowR * ((size_t)colsA * (size_t)2) + ((size_t)i * (size_t)2) + (size_t)1];
+        const float rB = ptrB[(size_t)i * ((size_t)colsB * (size_t)2) + ((size_t)colR * (size_t)2)];
+        const float iB = ptrB[(size_t)i * ((size_t)colsB * (size_t)2) + ((size_t)colR * (size_t)2) + (size_t)1];
+
+        //(rA + iA)(rB + iB)
+        const float first = rA * rB;
+        const float outer = rA * iB;
+        const float inner = iA * rB;
+        const float lasts = iA * iB;
+
+        rR += first + lasts;
+        iR += outer + inner;
+    }
+    ptrR[(size_t)rowR * ((size_t)colsR * (size_t)2) + ((size_t)colR * (size_t)2)] = rR;
+    ptrR[(size_t)rowR * ((size_t)colsR * (size_t)2) + ((size_t)colR * (size_t)2) + (size_t)1] = iR;
+}
+
+__kernel void kernel_knk
+(
+    __global float* ptrR,
+    __global float* ptrA,
+    __global float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB
+)
+{
+    const int rowsR = rowsA * rowsB;
+    const int colsR = colsA * colsB;
+    const int rowR = get_global_id(0); //{gpu_only}
+    for (int colR = 0; colR < colsR; colR++)
+    {
+        const int rowA = rowR / rowsB;
+        const int colA = colR / colsB;
+        const int rowB = rowR % rowsB;
+        const int colB = colR % colsB;
+
+        const int posA = rowA * (colsA * 2) + (colA * 2);
+        const int posB = rowB * (colsB * 2) + (colB * 2);
+
+        const float rA = ptrA[posA];
+        const float iA = ptrA[posA + 1];
+        const float rB = ptrB[posB];
+        const float iB = ptrB[posB + 1];
+
+        //(rA + iA)(rB + iB)
+        const float first = rA * rB;
+        const float outer = rA * iB;
+        const float inner = iA * rB;
+        const float lasts = iA * iB;
+        ptrR[rowR * (colsR * 2) + (colR * 2)] = first + lasts;
+        ptrR[rowR * (colsR * 2) + (colR * 2) + 1] = outer + inner;
+    }
+}
+
+__kernel void kernel_knk_2x2
+(
+    __global float* ptrR,
+    __global float* ptrA,
+    const int rowsA,
+    const int colsA,
+    const float gate0,
+    const float gate1,
+    const float gate2,
+    const float gate3,
+    const float gate4,
+    const float gate5,
+    const float gate6,
+    const float gate7
+)
+{
+    const int rowsR = rowsA * 2;
+    const int colsR = colsA * 2;
+    const int rowR = get_global_id(0) * 2; //{gpu_only}
+
+    for (int colR = 0; colR < colsR; colR += 2)
+    {
+        const int rowA = rowR / 2;
+        const int colA = colR / 2;
+        const float rA = ptrA[(size_t)rowA * ((size_t)colsA * (size_t)2) + ((size_t)colA * (size_t)2)];
+        const float iA = ptrA[(size_t)rowA * ((size_t)colsA * (size_t)2) + ((size_t)colA * (size_t)2) + (size_t)1];
+
+        for (int i = 0; i < 4; i++)
+        {
+            float rB, iB;
+            switch (i)
+            {
+                case 0: rB = gate0; iB = gate1; break;
+                case 1: rB = gate2; iB = gate3; break;
+                case 2: rB = gate4; iB = gate5; break;
+                case 3: rB = gate6; iB = gate7; break;
+            }
+
+            //(rA + iA)(rB + iB)
+            const float first = rA * rB;
+            const float outer = rA * iB;
+            const float inner = iA * rB;
+            const float lasts = iA * iB;
+            switch (i)
+            {
+                case 0:
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 1:
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 2:
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 3:
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+            }
+
+        }
+    }
+}
diff --git a/src/.kernel.tmp.2 b/src/.kernel.tmp.2
new file mode 100644 (file)
index 0000000..c0f9568
Binary files /dev/null and b/src/.kernel.tmp.2 differ
index 0ca140fee74c221c5f0d50a50c50a02e081cab2c..6577ebc159a691c9792f8a9ded6b99e861d7eb59 100644 (file)
@@ -1,9 +1,4 @@
-typedef struct
-{
-       char n[128];
-       unsigned char q0, q1, q2;
-       float arg0, arg1, arg2;
-} QInstr;
+#include "bytecode.h"
 
 const char* qansel_instruction_to_string(unsigned char instr)
 {
diff --git a/src/bytecode.h b/src/bytecode.h
new file mode 100644 (file)
index 0000000..f562277
--- /dev/null
@@ -0,0 +1,145 @@
+#ifndef __BYTECODE_H__
+#define __BYTECODE_H__
+
+#include <time.h>
+#include "complex.h"
+#include "context.h"
+#include "display.h"
+
+typedef struct
+{
+       char n[128];
+       unsigned char q0, q1, q2;
+       float arg0, arg1, arg2;
+} QInstr;
+
+typedef struct
+{
+       int size;
+       unsigned char bytes[16];
+       int barrier_width;
+       unsigned char barrier[QANSEL_QUBITS_MAX * 2];
+       int next;
+       int prev;
+       int checked;
+       int op;
+       unsigned char ifop[16];
+       int use_ifop;
+} QBytecode;
+
+static float Identity[] = 
+{
+    1, 0,    0, 0,
+    0, 0,    1, 0,
+};
+
+static float PauliX[] = 
+{
+    0, 0,    1, 0,
+    1, 0,    0, 0,
+};
+
+static float PauliY[] =
+{
+    0, 0,    0, 1,
+    0,-1,    0, 0,
+};
+
+static float PauliZ[] = 
+{
+    1, 0,    0, 0,
+    0, 0,   -1, 0,    
+};
+
+//  1/sqrt(2)
+#define R 0.7071067811865475
+static float Hadamard[] =
+{
+    R, 0,   R, 0,
+    R, 0,  -R, 0,
+};
+
+static float PhaseS[] =
+{
+    1, 0,    0, 0,
+    0, 0,    0,-1,
+};
+
+// 1/sqrt(2) + 1/sqrt(2)i
+static float PhaseT[] =
+{
+    1, 0,    0, 0,
+    0, 0,    R,-R,
+};
+
+static float ControlledNOT[] = 
+{
+    1, 0,    0, 0,    0, 0,    0, 0,
+    0, 0,    1, 0,    0, 0,    0, 0,   
+    0, 0,    0, 0,    0, 0,    1, 0,
+    0, 0,    0, 0,    1, 0,    0, 0,   
+};
+
+const char* qansel_instruction_to_string(unsigned char instr);
+void qansel_rand_init();
+float qansel_rand_s(float s);
+float qansel_rand_h();
+float qansel_rand_t(QAnselContext* ctx);
+void qansel_cnot(cpx_mtx_t* stateVector, unsigned char qubitCount, unsigned char bitA, unsigned char bitB);
+void qansel_fredkin(cpx_mtx_t* stateVector, unsigned char qubitCount, unsigned char bitA, unsigned char bitB, unsigned char bitC);
+void qansel_toffoli(cpx_mtx_t* stateVector, unsigned char qubitCount, unsigned char bitA, unsigned char bitB, unsigned char bitC);
+float* qansel_unitary(float theta, float phi, float lambda);
+
+void qansel_queue_init
+(
+       cpx_mtx_t* queueVector,
+       int qubitCount,
+       int dofree
+);
+
+void qansel_instruction_queue
+(
+       cpx_mtx_t** queueVector,
+       int qubitCount,
+       unsigned char instr,
+       unsigned char index, 
+       float arg0,
+       float arg1,
+       float arg2
+);
+
+void qansel_instruction
+(
+       QAnselContext* ctx,
+       cpx_mtx_t* stateVector,
+       int qubitCount,
+       unsigned char instr,
+       unsigned char index, 
+       float arg0,
+       float arg1,
+       float arg2,
+       //pass in this to flush the queue
+       cpx_mtx_t** queueVector
+);
+unsigned char qansel_measure(QAnselContext* ctx, cpx_mtx_t* stateVector, unsigned char qubitCount, unsigned char qubit);
+int qansel_get_instruction_bitmax(unsigned char* ptr, int offset, int* bitmax, int* qbitmax);
+int qansel_get_instruction_size(unsigned char instr);
+void qansel_get_barrier(QBytecode** qbc, int idx);
+void qansel_born(cpx_mtx_t* stateVector, int PC, int qubitCount, unsigned char q0);
+void qansel_density_or_print(cpx_mtx_t* stateVector, unsigned char* bitVector, unsigned char density, int bitCount, int qubitCount, unsigned char a0);
+float qansel_get_float(unsigned char* program, int offset);
+short qansel_get_short(unsigned char* program, int offset);
+int qansel_get_int(unsigned char* program, int offset);
+void qansel_reset(QAnselContext* ctx, cpx_mtx_t* stateVector, unsigned char* bitVector, int qubitCount, int bitCount, unsigned char q0);
+unsigned char qansel_compare(unsigned char* bitVector, int bitCount, int PC, unsigned char a0, short op);
+
+//computes program efficiency
+//     points are awarded for how long strings of single qubit
+//     instructions are 
+int qansel_efficiency(QBytecode* program, int programSize, int head);
+int qansel_crawl(QAnselContext* ctx, unsigned char* program, int programSize, int* qubitCount, int* bitCount);
+void qansel_reorder(QAnselContext* ctx, unsigned char* program, int programSize);
+void qansel_run(QAnselContext* ctx, unsigned char* program, int programSize, int qubitCount, int bitCount, unsigned char* outputBitVector);
+int qanselExecuteBytecode(unsigned char* buff, int sizeofbuff, QAnselContext* ctx);
+
+#endif
\ No newline at end of file
diff --git a/src/bytecode.o b/src/bytecode.o
new file mode 100644 (file)
index 0000000..4379a3f
Binary files /dev/null and b/src/bytecode.o differ
index 6abeeb546277d9c0ff99ec71ad3013b1944cf684..8916788c1dc408d41d4cc59d97494784f727b8af 100644 (file)
@@ -1,16 +1,4 @@
-#ifndef __cpx__
-#define __cpx__
-typedef struct
-{
-    float real, imaginary;
-} cpx_t;
-
-typedef struct
-{
-    float *ptr;
-    int rows, cols;
-} cpx_mtx_t;
-
+#include "complex.h"
 unsigned char* cpx_str(cpx_t* n)
 {
     unsigned char* r;
@@ -183,18 +171,6 @@ void cpx_mtx_knk_2x2(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA
 /*-----------------------------------------------------------------------------------*/
 /*THREADED*/
 /*-----------------------------------------------------------------------------------*/
-typedef struct
-{
-       float* ptrR;
-       float* ptrA;
-       float* ptrB;
-       int rowsA;
-       int colsA;
-       int rowsB;
-       int colsB;
-       int delimeterStart;
-       int delimeterCount;
-} cpx_thread_context;
 
 void* cpx_mtx_knk_threads_run(void *context)
 {
@@ -822,6 +798,3 @@ void cpx_mtx_knk_metal_2x2(float* ptrR, float* ptrA, float* ptrB, int rowsA, int
        cpx_mtx_knk_metal_2x2(ptrR, ptrA, ptrB, rowsA, colsA, rowsB, colsB);
        #endif
 }
-
-
-#endif
\ No newline at end of file
diff --git a/src/complex.h b/src/complex.h
new file mode 100644 (file)
index 0000000..41efbac
--- /dev/null
@@ -0,0 +1,101 @@
+#ifndef __COMPLEX_H__
+#define __COMPLEX_H__
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <stddef.h>
+#include <math.h>
+#include "hardware.h"
+#include "kernel.h"
+
+#ifdef __PTHREAD__
+#include <pthread.h>
+#endif
+
+#ifdef __OPENCL__
+#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
+#define CL_TARGET_OPENCL_VERSION 300
+#include <CL/cl.h>
+#define gpuerr(x) if (err != CL_SUCCESS) { fprintf(stderr, "GPU Error on line %i: %s.\n", __LINE__, clGetErrorString(x)); exit(1); }
+static cl_platform_id cpx_mtx_platform_id;
+static cl_device_id cpx_mtx_device_id;
+static cl_context cpx_mtx_context;
+static cl_command_queue cpx_mtx_command_queue;
+static unsigned char* cpx_mtx_cache = NULL;
+static size_t cpx_mtx_cache_len = 0;
+#include "kernel_gpu.cl"
+#endif
+
+typedef struct
+{
+    float real, imaginary;
+} cpx_t;
+
+typedef struct
+{
+    float *ptr;
+    int rows, cols;
+} cpx_mtx_t;
+
+unsigned char* cpx_str(cpx_t* n);
+cpx_t cpx_new(float r, float i);
+void cpx_mul(cpx_t* r, cpx_t* a, cpx_t* b);
+float cpx_magsqr(cpx_t* n);
+float cpx_mag(cpx_t* n);
+void cpx_mtx_set(cpx_mtx_t* m, int row, int col, cpx_t* n);
+void cpx_mtx_set2(cpx_mtx_t* m, int row, int col, float real, float imaginary);
+void cpx_mtx_get(cpx_mtx_t* m, int row, int col, cpx_t* n);
+void cpx_mtx_copy(cpx_mtx_t* m, cpx_mtx_t* n);
+int cpx_mtx_cmp(cpx_mtx_t* m, cpx_mtx_t* n);
+void cpx_mtx_init(cpx_mtx_t* m, int rows, int cols);
+void cpx_mtx_free(cpx_mtx_t* m);
+void cpx_mtx_print(cpx_mtx_t* m);
+void cpx_mtx_dot(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+void cpx_mtx_knk(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+void cpx_mtx_knk_2x2(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+/*-----------------------------------------------------------------------------------*/
+/*THREADED*/
+/*-----------------------------------------------------------------------------------*/
+typedef struct
+{
+       float* ptrR;
+       float* ptrA;
+       float* ptrB;
+       int rowsA;
+       int colsA;
+       int rowsB;
+       int colsB;
+       int delimeterStart;
+       int delimeterCount;
+} cpx_thread_context;
+
+void* cpx_mtx_knk_threads_run(void *context);
+void cpx_mtx_knk_threads(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+void* cpx_mtx_knk_threads_2x2_run(void *context);
+
+void cpx_mtx_knk_threads_2x2(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+void* cpx_mtx_dot_threads_run(void *context);
+void cpx_mtx_dot_threads(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+
+/*-----------------------------------------------------------------------------------*/
+/*METAL*/
+/*-----------------------------------------------------------------------------------*/
+#ifdef __OPENCL__
+const char* clGetErrorString(cl_int err);
+#endif
+
+unsigned char cpx_mtx_begin(unsigned char verbose);
+
+void cpx_mtx_clean();
+
+void cpx_mtx_dot_metal(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+
+void cpx_mtx_knk_metal(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+
+//This only works if ptrA is NxM where both N and X are divisible by two,
+//     and ptrB is 2x2. If both are true, this is much more efficient than
+//     the standard knk_metal() function. 
+void cpx_mtx_knk_metal_2x2(float* ptrR, float* ptrA, float* ptrB, int rowsA, int colsA, int rowsB, int colsB);
+
+#endif
\ No newline at end of file
diff --git a/src/complex.o b/src/complex.o
new file mode 100644 (file)
index 0000000..7183f02
Binary files /dev/null and b/src/complex.o differ
index e25abc8de7cfedf5114bfaf604508e1f01edd1f9..04503bcc657888b2fccd0662e2d19968941ba66e 100644 (file)
@@ -1,3 +1,4 @@
+#include "context.h"
 int qanselContextValidate(QAnselContext* ctx)
 {
        if (ctx->memory_limit == 0) ctx->memory_limit = QANSEL_QUBITS_MAX;
diff --git a/src/context.h b/src/context.h
new file mode 100644 (file)
index 0000000..a80670f
--- /dev/null
@@ -0,0 +1,26 @@
+#ifndef __CONTEXT_H__
+#define __CONTEXT_H__
+
+#include <stdio.h>
+#include "qansel.h"
+#include "complex.h"
+#include "hardware.h"
+
+typedef struct
+{
+       unsigned char memory_limit;
+       unsigned char hardware_rng;
+       unsigned char verbose;
+       unsigned char optimization_level;
+       unsigned char sampling_bit;
+       float hidden_variable;
+       int display_delay;
+       int sampling_shots;
+       FILE* random_file;
+} QAnselContext;
+
+int qanselContextValidate(QAnselContext* ctx);
+int qanselContextBegin(QAnselContext* ctx);
+int qanselContextEnd(QAnselContext* ctx);
+
+#endif
\ No newline at end of file
diff --git a/src/context.o b/src/context.o
new file mode 100644 (file)
index 0000000..83f557e
Binary files /dev/null and b/src/context.o differ
index 489d3788581c8eee238043811941078592290fca..b0519b6b1016ff3b5b06bb5cad3112678576af6d 100644 (file)
@@ -1,3 +1,7 @@
+#include "qansel.h"
+#include "complex.h"
+#include "display.h"
+
 #ifdef __SDL2__
 void DrawThickLine(SDL_Renderer* renderer, int x1, int y1, int x2, int y2, int thickness) {
     int deltaX = x2 - x1;
diff --git a/src/display.h b/src/display.h
new file mode 100644 (file)
index 0000000..9c555ad
--- /dev/null
@@ -0,0 +1,11 @@
+#ifndef __DISPLAY_H__
+#define __DISPLAY_H__
+
+#include <stdint.h>
+#ifdef __SDL2__
+#include <SDL.h>
+void DrawThickLine(SDL_Renderer* renderer, int x1, int y1, int x2, int y2, int thickness);
+#endif
+int display(cpx_mtx_t* stateVector, uint8_t qubitCount, int delay);
+
+#endif
\ No newline at end of file
diff --git a/src/display.o b/src/display.o
new file mode 100644 (file)
index 0000000..8e7c04b
Binary files /dev/null and b/src/display.o differ
diff --git a/src/finaldone b/src/finaldone
new file mode 100755 (executable)
index 0000000..925897b
Binary files /dev/null and b/src/finaldone differ
diff --git a/src/gates.c b/src/gates.c
deleted file mode 100644 (file)
index 63bbf90..0000000
+++ /dev/null
@@ -1,53 +0,0 @@
-
-float Identity[] = 
-{
-    1, 0,    0, 0,
-    0, 0,    1, 0,
-};
-
-float PauliX[] = 
-{
-    0, 0,    1, 0,
-    1, 0,    0, 0,
-};
-
-float PauliY[] =
-{
-    0, 0,    0, 1,
-    0,-1,    0, 0,
-};
-
-float PauliZ[] = 
-{
-    1, 0,    0, 0,
-    0, 0,   -1, 0,    
-};
-
-//  1/sqrt(2)
-#define R 0.7071067811865475
-float Hadamard[] =
-{
-    R, 0,   R, 0,
-    R, 0,  -R, 0,
-};
-
-float PhaseS[] =
-{
-    1, 0,    0, 0,
-    0, 0,    0,-1,
-};
-
-// 1/sqrt(2) + 1/sqrt(2)i
-float PhaseT[] =
-{
-    1, 0,    0, 0,
-    0, 0,    R,-R,
-};
-
-float ControlledNOT[] = 
-{
-    1, 0,    0, 0,    0, 0,    0, 0,
-    0, 0,    1, 0,    0, 0,    0, 0,   
-    0, 0,    0, 0,    0, 0,    1, 0,
-    0, 0,    0, 0,    1, 0,    0, 0,   
-};
\ No newline at end of file
index 130d87e1814a95bcb438b2e6cc278d4a03c6b94b..0283ca2a48bdb4ebaf92a676256a7352bff1f567 100644 (file)
@@ -1,3 +1,5 @@
+#include "hardware.h"
+
 int qansel___get_core_count()
 {
     #if defined(_WIN32) || defined(_WIN64)
diff --git a/src/hardware.h b/src/hardware.h
new file mode 100644 (file)
index 0000000..d0cf307
--- /dev/null
@@ -0,0 +1,20 @@
+#ifndef __HARDWARE_H__
+#define __HARDWARE_H__
+
+#include <stddef.h>
+#if defined(_WIN32) || defined(_WIN64)
+#include <windows.h>
+#elif defined(__linux__)
+#include <unistd.h>
+#include <sys/time.h>
+#elif defined(__APPLE__)
+#include <sys/sysctl.h>
+#endif
+
+int qansel___get_core_count();
+int qansel_get_core_count();
+unsigned long int qansel_get_time();
+int qansel_hardware_rand_supported();
+unsigned char qansel_hardware_rand();
+
+#endif
\ No newline at end of file
diff --git a/src/hardware.o b/src/hardware.o
new file mode 100644 (file)
index 0000000..3a95cd2
Binary files /dev/null and b/src/hardware.o differ
diff --git a/src/interface.html b/src/interface.html
deleted file mode 100644 (file)
index de7e2e4..0000000
+++ /dev/null
@@ -1,40 +0,0 @@
-<body>
-<table>
-       <tr>
-               <td>H</td>
-               <td>X</td>
-               <td>Y</td>
-               <td>Z</td>
-       </tr>
-</table>
-<table>
-       <tr>
-               <td>Q0</td>
-               <td>-</td>
-               <td>-</td>
-               <td>-</td>
-       </tr>
-       <tr>
-               <td>Q1</td>
-               <td>-</td>
-               <td>-</td>
-               <td>-</td>
-       </tr>
-       <tr>
-               <td>Q2</td>
-       </tr>
-       <tr>
-               <td>Q3</td>
-               <td>-</td>
-               <td>-</td>
-               <td>-</td>
-       </tr>
-       <tr>
-               <td>Q4</td>
-               <td>-</td>
-               <td>-</td>
-               <td>-</td>
-       </tr>
-</table>
-
-</body>
\ No newline at end of file
index dcfdf50a659f6d8d9f1ec1e4dcc2783b8822364d..a5731c3617384a2cea8a3bcb4718d4b022169ec5 100644 (file)
@@ -1,3 +1,5 @@
+#include "kernel.h" //{cpu_only}}
+
 __kernel void kernel_dot
 (
     __global float* ptrR,
diff --git a/src/kernel.h b/src/kernel.h
new file mode 100644 (file)
index 0000000..f71b799
--- /dev/null
@@ -0,0 +1,48 @@
+#ifndef __KERNEL_H__
+#define __KERNEL_H__
+
+#include <stddef.h>
+
+void kernel_dot
+(
+    float* ptrR,
+    float* ptrA,
+    float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB,
+    const int get_global_id_0,
+    const int get_global_id_1
+);
+
+void kernel_knk
+(
+    float* ptrR,
+    float* ptrA,
+    float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB,
+    const int get_global_id_0
+);
+
+void kernel_knk_2x2
+(
+    float* ptrR,
+    float* ptrA,
+    const int rowsA,
+    const int colsA,
+    const float gate0,
+    const float gate1,
+    const float gate2,
+    const float gate3,
+    const float gate4,
+    const float gate5,
+    const float gate6,
+    const float gate7,
+       const int get_global_id_0
+);
+
+#endif
\ No newline at end of file
diff --git a/src/kernel_cpu.c b/src/kernel_cpu.c
new file mode 100644 (file)
index 0000000..842f3c5
--- /dev/null
@@ -0,0 +1,150 @@
+#include "kernel.h" //{cpu_only}}
+
+void kernel_dot
+(
+    float* ptrR,
+    float* ptrA,
+    float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB
+    ,   const int get_global_id_0, //{cpu_only}
+        const int get_global_id_1 //{cpu_only}
+)
+{
+    const int rowsR = rowsA;
+    const int colsR = colsB;
+    const int rowR = get_global_id_0; //{cpu_only}
+    const int colR = get_global_id_1; //{cpu_only}
+    
+    float rR = 0;
+    float iR = 0;
+
+    for (int i = 0; i < colsA; i++)
+    {
+        const float rA = ptrA[(size_t)rowR * ((size_t)colsA * (size_t)2) + ((size_t)i * (size_t)2)];
+        const float iA = ptrA[(size_t)rowR * ((size_t)colsA * (size_t)2) + ((size_t)i * (size_t)2) + (size_t)1];
+        const float rB = ptrB[(size_t)i * ((size_t)colsB * (size_t)2) + ((size_t)colR * (size_t)2)];
+        const float iB = ptrB[(size_t)i * ((size_t)colsB * (size_t)2) + ((size_t)colR * (size_t)2) + (size_t)1];
+
+        //(rA + iA)(rB + iB)
+        const float first = rA * rB;
+        const float outer = rA * iB;
+        const float inner = iA * rB;
+        const float lasts = iA * iB;
+
+        rR += first + lasts;
+        iR += outer + inner;
+    }
+    ptrR[(size_t)rowR * ((size_t)colsR * (size_t)2) + ((size_t)colR * (size_t)2)] = rR;
+    ptrR[(size_t)rowR * ((size_t)colsR * (size_t)2) + ((size_t)colR * (size_t)2) + (size_t)1] = iR;
+}
+
+void kernel_knk
+(
+    float* ptrR,
+    float* ptrA,
+    float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB
+    ,   const int get_global_id_0 //{cpu_only}
+)
+{
+    const int rowsR = rowsA * rowsB;
+    const int colsR = colsA * colsB;
+    const int rowR = get_global_id_0; //{cpu_only}
+    for (int colR = 0; colR < colsR; colR++)
+    {
+        const int rowA = rowR / rowsB;
+        const int colA = colR / colsB;
+        const int rowB = rowR % rowsB;
+        const int colB = colR % colsB;
+
+        const int posA = rowA * (colsA * 2) + (colA * 2);
+        const int posB = rowB * (colsB * 2) + (colB * 2);
+
+        const float rA = ptrA[posA];
+        const float iA = ptrA[posA + 1];
+        const float rB = ptrB[posB];
+        const float iB = ptrB[posB + 1];
+
+        //(rA + iA)(rB + iB)
+        const float first = rA * rB;
+        const float outer = rA * iB;
+        const float inner = iA * rB;
+        const float lasts = iA * iB;
+        ptrR[rowR * (colsR * 2) + (colR * 2)] = first + lasts;
+        ptrR[rowR * (colsR * 2) + (colR * 2) + 1] = outer + inner;
+    }
+}
+
+void kernel_knk_2x2
+(
+    float* ptrR,
+    float* ptrA,
+    const int rowsA,
+    const int colsA,
+    const float gate0,
+    const float gate1,
+    const float gate2,
+    const float gate3,
+    const float gate4,
+    const float gate5,
+    const float gate6,
+    const float gate7
+    , const int get_global_id_0 //{cpu_only}
+)
+{
+    const int rowsR = rowsA * 2;
+    const int colsR = colsA * 2;
+    const int rowR = get_global_id_0 * 2; //{cpu_only}
+
+    for (int colR = 0; colR < colsR; colR += 2)
+    {
+        const int rowA = rowR / 2;
+        const int colA = colR / 2;
+        const float rA = ptrA[(size_t)rowA * ((size_t)colsA * (size_t)2) + ((size_t)colA * (size_t)2)];
+        const float iA = ptrA[(size_t)rowA * ((size_t)colsA * (size_t)2) + ((size_t)colA * (size_t)2) + (size_t)1];
+
+        for (int i = 0; i < 4; i++)
+        {
+            float rB, iB;
+            switch (i)
+            {
+                case 0: rB = gate0; iB = gate1; break;
+                case 1: rB = gate2; iB = gate3; break;
+                case 2: rB = gate4; iB = gate5; break;
+                case 3: rB = gate6; iB = gate7; break;
+            }
+
+            //(rA + iA)(rB + iB)
+            const float first = rA * rB;
+            const float outer = rA * iB;
+            const float inner = iA * rB;
+            const float lasts = iA * iB;
+            switch (i)
+            {
+                case 0:
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 1:
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 2:
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 3:
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+            }
+
+        }
+    }
+}
diff --git a/src/kernel_cpu.cl b/src/kernel_cpu.cl
new file mode 100644 (file)
index 0000000..842f3c5
--- /dev/null
@@ -0,0 +1,150 @@
+#include "kernel.h" //{cpu_only}}
+
+void kernel_dot
+(
+    float* ptrR,
+    float* ptrA,
+    float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB
+    ,   const int get_global_id_0, //{cpu_only}
+        const int get_global_id_1 //{cpu_only}
+)
+{
+    const int rowsR = rowsA;
+    const int colsR = colsB;
+    const int rowR = get_global_id_0; //{cpu_only}
+    const int colR = get_global_id_1; //{cpu_only}
+    
+    float rR = 0;
+    float iR = 0;
+
+    for (int i = 0; i < colsA; i++)
+    {
+        const float rA = ptrA[(size_t)rowR * ((size_t)colsA * (size_t)2) + ((size_t)i * (size_t)2)];
+        const float iA = ptrA[(size_t)rowR * ((size_t)colsA * (size_t)2) + ((size_t)i * (size_t)2) + (size_t)1];
+        const float rB = ptrB[(size_t)i * ((size_t)colsB * (size_t)2) + ((size_t)colR * (size_t)2)];
+        const float iB = ptrB[(size_t)i * ((size_t)colsB * (size_t)2) + ((size_t)colR * (size_t)2) + (size_t)1];
+
+        //(rA + iA)(rB + iB)
+        const float first = rA * rB;
+        const float outer = rA * iB;
+        const float inner = iA * rB;
+        const float lasts = iA * iB;
+
+        rR += first + lasts;
+        iR += outer + inner;
+    }
+    ptrR[(size_t)rowR * ((size_t)colsR * (size_t)2) + ((size_t)colR * (size_t)2)] = rR;
+    ptrR[(size_t)rowR * ((size_t)colsR * (size_t)2) + ((size_t)colR * (size_t)2) + (size_t)1] = iR;
+}
+
+void kernel_knk
+(
+    float* ptrR,
+    float* ptrA,
+    float* ptrB,
+    const int rowsA,
+    const int colsA,
+    const int rowsB,
+    const int colsB
+    ,   const int get_global_id_0 //{cpu_only}
+)
+{
+    const int rowsR = rowsA * rowsB;
+    const int colsR = colsA * colsB;
+    const int rowR = get_global_id_0; //{cpu_only}
+    for (int colR = 0; colR < colsR; colR++)
+    {
+        const int rowA = rowR / rowsB;
+        const int colA = colR / colsB;
+        const int rowB = rowR % rowsB;
+        const int colB = colR % colsB;
+
+        const int posA = rowA * (colsA * 2) + (colA * 2);
+        const int posB = rowB * (colsB * 2) + (colB * 2);
+
+        const float rA = ptrA[posA];
+        const float iA = ptrA[posA + 1];
+        const float rB = ptrB[posB];
+        const float iB = ptrB[posB + 1];
+
+        //(rA + iA)(rB + iB)
+        const float first = rA * rB;
+        const float outer = rA * iB;
+        const float inner = iA * rB;
+        const float lasts = iA * iB;
+        ptrR[rowR * (colsR * 2) + (colR * 2)] = first + lasts;
+        ptrR[rowR * (colsR * 2) + (colR * 2) + 1] = outer + inner;
+    }
+}
+
+void kernel_knk_2x2
+(
+    float* ptrR,
+    float* ptrA,
+    const int rowsA,
+    const int colsA,
+    const float gate0,
+    const float gate1,
+    const float gate2,
+    const float gate3,
+    const float gate4,
+    const float gate5,
+    const float gate6,
+    const float gate7
+    , const int get_global_id_0 //{cpu_only}
+)
+{
+    const int rowsR = rowsA * 2;
+    const int colsR = colsA * 2;
+    const int rowR = get_global_id_0 * 2; //{cpu_only}
+
+    for (int colR = 0; colR < colsR; colR += 2)
+    {
+        const int rowA = rowR / 2;
+        const int colA = colR / 2;
+        const float rA = ptrA[(size_t)rowA * ((size_t)colsA * (size_t)2) + ((size_t)colA * (size_t)2)];
+        const float iA = ptrA[(size_t)rowA * ((size_t)colsA * (size_t)2) + ((size_t)colA * (size_t)2) + (size_t)1];
+
+        for (int i = 0; i < 4; i++)
+        {
+            float rB, iB;
+            switch (i)
+            {
+                case 0: rB = gate0; iB = gate1; break;
+                case 1: rB = gate2; iB = gate3; break;
+                case 2: rB = gate4; iB = gate5; break;
+                case 3: rB = gate6; iB = gate7; break;
+            }
+
+            //(rA + iA)(rB + iB)
+            const float first = rA * rB;
+            const float outer = rA * iB;
+            const float inner = iA * rB;
+            const float lasts = iA * iB;
+            switch (i)
+            {
+                case 0:
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 1:
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)0) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 2:
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)0) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+                case 3:
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2)] = first + lasts;
+                    ptrR[((size_t)rowR + (size_t)1) * ((size_t)colsR * (size_t)2) + (((size_t)colR + (size_t)1) * (size_t)2) + (size_t)1] = outer + inner;
+                break;
+            }
+
+        }
+    }
+}
diff --git a/src/kernel_cpu.o b/src/kernel_cpu.o
new file mode 100644 (file)
index 0000000..baa933a
Binary files /dev/null and b/src/kernel_cpu.o differ
diff --git a/src/kernel_gpu.c b/src/kernel_gpu.c
new file mode 100644 (file)
index 0000000..c489f41
--- /dev/null
@@ -0,0 +1,436 @@
+static unsigned char kernel_gpu[] = {
+  0x0a, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x20, 0x76, 0x6f,
+  0x69, 0x64, 0x20, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x64, 0x6f,
+  0x74, 0x0a, 0x28, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c,
+  0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20,
+  0x70, 0x74, 0x72, 0x52, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f,
+  0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x2a, 0x20, 0x70, 0x74, 0x72, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x42, 0x2c, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20,
+  0x72, 0x6f, 0x77, 0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c,
+  0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x42, 0x2c,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69,
+  0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x0a, 0x29, 0x0a, 0x7b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69,
+  0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x52, 0x20, 0x3d, 0x20, 0x72,
+  0x6f, 0x77, 0x73, 0x41, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73,
+  0x52, 0x20, 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f,
+  0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29,
+  0x3b, 0x20, 0x2f, 0x2f, 0x7b, 0x67, 0x70, 0x75, 0x5f, 0x6f, 0x6e, 0x6c,
+  0x79, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x3d, 0x20,
+  0x67, 0x65, 0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69,
+  0x64, 0x28, 0x31, 0x29, 0x3b, 0x20, 0x2f, 0x2f, 0x7b, 0x67, 0x70, 0x75,
+  0x5f, 0x6f, 0x6e, 0x6c, 0x79, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x52,
+  0x20, 0x3d, 0x20, 0x30, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x69, 0x52, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x0a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e,
+  0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c,
+  0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x3b, 0x20, 0x69, 0x2b, 0x2b, 0x29,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x72, 0x41, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x41,
+  0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77,
+  0x52, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x69, 0x20, 0x2a, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x41, 0x20, 0x3d,
+  0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x69, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x31, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x42, 0x5b, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x69, 0x20, 0x2a, 0x20, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73,
+  0x42, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x70,
+  0x74, 0x72, 0x42, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x69, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31,
+  0x5d, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x2f, 0x2f, 0x28, 0x72, 0x41, 0x20, 0x2b, 0x20, 0x69, 0x41, 0x29, 0x28,
+  0x72, 0x42, 0x20, 0x2b, 0x20, 0x69, 0x42, 0x29, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x3d,
+  0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20,
+  0x3d, 0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72,
+  0x20, 0x3d, 0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6c, 0x61, 0x73, 0x74,
+  0x73, 0x20, 0x3d, 0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b,
+  0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x72, 0x52,
+  0x20, 0x2b, 0x3d, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x2b, 0x20,
+  0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x69, 0x52, 0x20, 0x2b, 0x3d, 0x20, 0x6f, 0x75, 0x74,
+  0x65, 0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74,
+  0x72, 0x52, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72,
+  0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20,
+  0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x52, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x69, 0x52,
+  0x3b, 0x0a, 0x7d, 0x0a, 0x0a, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x5f, 0x6b, 0x6e, 0x6b, 0x0a, 0x28, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x52, 0x2c, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x41, 0x2c, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x42,
+  0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x2c, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f,
+  0x77, 0x73, 0x42, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42,
+  0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x52,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x72,
+  0x6f, 0x77, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73,
+  0x52, 0x20, 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20,
+  0x63, 0x6f, 0x6c, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77,
+  0x52, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62,
+  0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29, 0x3b, 0x20, 0x2f, 0x2f,
+  0x7b, 0x67, 0x70, 0x75, 0x5f, 0x6f, 0x6e, 0x6c, 0x79, 0x7d, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20,
+  0x63, 0x6f, 0x6c, 0x52, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x63, 0x6f,
+  0x6c, 0x52, 0x20, 0x3c, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x3b, 0x20,
+  0x63, 0x6f, 0x6c, 0x52, 0x2b, 0x2b, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x41,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2f, 0x20, 0x72, 0x6f,
+  0x77, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63,
+  0x6f, 0x6c, 0x41, 0x20, 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2f,
+  0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e,
+  0x74, 0x20, 0x72, 0x6f, 0x77, 0x42, 0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77,
+  0x52, 0x20, 0x25, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x42, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x42, 0x20, 0x3d, 0x20,
+  0x63, 0x6f, 0x6c, 0x52, 0x20, 0x25, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42,
+  0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x70, 0x6f, 0x73,
+  0x41, 0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x41, 0x20, 0x2a, 0x20, 0x28,
+  0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b,
+  0x20, 0x28, 0x63, 0x6f, 0x6c, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x70, 0x6f, 0x73, 0x42, 0x20,
+  0x3d, 0x20, 0x72, 0x6f, 0x77, 0x42, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f,
+  0x6c, 0x73, 0x42, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x63, 0x6f, 0x6c, 0x42, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x3b, 0x0a, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x41, 0x20, 0x3d,
+  0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x70, 0x6f, 0x73, 0x41, 0x5d, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x41, 0x20,
+  0x3d, 0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x70, 0x6f, 0x73, 0x41, 0x20,
+  0x2b, 0x20, 0x31, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x42, 0x5b,
+  0x70, 0x6f, 0x73, 0x42, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x42,
+  0x5b, 0x70, 0x6f, 0x73, 0x42, 0x20, 0x2b, 0x20, 0x31, 0x5d, 0x3b, 0x0a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x2f, 0x2f, 0x28,
+  0x72, 0x41, 0x20, 0x2b, 0x20, 0x69, 0x41, 0x29, 0x28, 0x72, 0x42, 0x20,
+  0x2b, 0x20, 0x69, 0x42, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x3d, 0x20, 0x72, 0x41,
+  0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x3d, 0x20, 0x72,
+  0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x20, 0x3d, 0x20,
+  0x69, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x20, 0x3d,
+  0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x72,
+  0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x52,
+  0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69,
+  0x72, 0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72,
+  0x52, 0x5b, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f,
+  0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20,
+  0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x2b,
+  0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7d, 0x0a, 0x7d, 0x0a, 0x0a, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x5f, 0x6b, 0x6e, 0x6b, 0x5f, 0x32, 0x78, 0x32, 0x0a, 0x28, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x52,
+  0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62,
+  0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74,
+  0x72, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x2c,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69,
+  0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x20, 0x67, 0x61, 0x74, 0x65, 0x30, 0x2c, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x67, 0x61, 0x74, 0x65, 0x31, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
+  0x67, 0x61, 0x74, 0x65, 0x32, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67,
+  0x61, 0x74, 0x65, 0x33, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61,
+  0x74, 0x65, 0x34, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61, 0x74,
+  0x65, 0x35, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61, 0x74, 0x65,
+  0x36, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61, 0x74, 0x65, 0x37,
+  0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x52,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x3d, 0x20,
+  0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f,
+  0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29,
+  0x20, 0x2a, 0x20, 0x32, 0x3b, 0x20, 0x2f, 0x2f, 0x7b, 0x67, 0x70, 0x75,
+  0x5f, 0x6f, 0x6e, 0x6c, 0x79, 0x7d, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20,
+  0x3c, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x3b, 0x20, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x2b, 0x3d, 0x20, 0x32, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x41,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2f, 0x20, 0x32, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x41, 0x20,
+  0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2f, 0x20, 0x32, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x41, 0x20, 0x3d,
+  0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x72, 0x6f, 0x77, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x63, 0x6f, 0x6c, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x69, 0x41, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72,
+  0x41, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f,
+  0x77, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x41,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x31, 0x5d, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20,
+  0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x34, 0x3b, 0x20,
+  0x69, 0x2b, 0x2b, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x42, 0x2c,
+  0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x73, 0x77, 0x69, 0x74, 0x63, 0x68, 0x20,
+  0x28, 0x69, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61,
+  0x73, 0x65, 0x20, 0x30, 0x3a, 0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x67,
+  0x61, 0x74, 0x65, 0x30, 0x3b, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x67,
+  0x61, 0x74, 0x65, 0x31, 0x3b, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20, 0x31, 0x3a,
+  0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65, 0x32, 0x3b,
+  0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65, 0x33, 0x3b,
+  0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x63, 0x61, 0x73, 0x65, 0x20, 0x32, 0x3a, 0x20, 0x72, 0x42, 0x20, 0x3d,
+  0x20, 0x67, 0x61, 0x74, 0x65, 0x34, 0x3b, 0x20, 0x69, 0x42, 0x20, 0x3d,
+  0x20, 0x67, 0x61, 0x74, 0x65, 0x35, 0x3b, 0x20, 0x62, 0x72, 0x65, 0x61,
+  0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20,
+  0x33, 0x3a, 0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65,
+  0x36, 0x3b, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65,
+  0x37, 0x3b, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x2f, 0x2f, 0x28, 0x72, 0x41, 0x20, 0x2b, 0x20, 0x69, 0x41, 0x29,
+  0x28, 0x72, 0x42, 0x20, 0x2b, 0x20, 0x69, 0x42, 0x29, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x66, 0x69,
+  0x72, 0x73, 0x74, 0x20, 0x3d, 0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x72,
+  0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x3d, 0x20, 0x72,
+  0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x6e, 0x6e, 0x65,
+  0x72, 0x20, 0x3d, 0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x20, 0x3d, 0x20, 0x69, 0x41, 0x20,
+  0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x73, 0x77, 0x69, 0x74, 0x63, 0x68,
+  0x20, 0x28, 0x69, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x61, 0x73, 0x65, 0x20, 0x30, 0x3a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a,
+  0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f,
+  0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a,
+  0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d,
+  0x20, 0x3d, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c,
+  0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20,
+  0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c,
+  0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b,
+  0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20,
+  0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e,
+  0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x62, 0x72, 0x65,
+  0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65,
+  0x20, 0x31, 0x3a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20,
+  0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74,
+  0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70,
+  0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f,
+  0x75, 0x74, 0x65, 0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20, 0x32, 0x3a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72,
+  0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72,
+  0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b,
+  0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63,
+  0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69, 0x72,
+  0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52,
+  0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f,
+  0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20,
+  0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f,
+  0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65,
+  0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20, 0x33, 0x3a, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52,
+  0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31,
+  0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52,
+  0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31,
+  0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20,
+  0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20,
+  0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x29,
+  0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20,
+  0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x29,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x2b,
+  0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7d, 0x0a, 0x7d, 0x00
+};
+static unsigned int kernel_gpu_len = 5188;
diff --git a/src/kernel_gpu.cl b/src/kernel_gpu.cl
new file mode 100644 (file)
index 0000000..c489f41
--- /dev/null
@@ -0,0 +1,436 @@
+static unsigned char kernel_gpu[] = {
+  0x0a, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x20, 0x76, 0x6f,
+  0x69, 0x64, 0x20, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x5f, 0x64, 0x6f,
+  0x74, 0x0a, 0x28, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c,
+  0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20,
+  0x70, 0x74, 0x72, 0x52, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f,
+  0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x2a, 0x20, 0x70, 0x74, 0x72, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x42, 0x2c, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20,
+  0x72, 0x6f, 0x77, 0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c,
+  0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x42, 0x2c,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69,
+  0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x0a, 0x29, 0x0a, 0x7b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69,
+  0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x52, 0x20, 0x3d, 0x20, 0x72,
+  0x6f, 0x77, 0x73, 0x41, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73,
+  0x52, 0x20, 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f,
+  0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29,
+  0x3b, 0x20, 0x2f, 0x2f, 0x7b, 0x67, 0x70, 0x75, 0x5f, 0x6f, 0x6e, 0x6c,
+  0x79, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x3d, 0x20,
+  0x67, 0x65, 0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69,
+  0x64, 0x28, 0x31, 0x29, 0x3b, 0x20, 0x2f, 0x2f, 0x7b, 0x67, 0x70, 0x75,
+  0x5f, 0x6f, 0x6e, 0x6c, 0x79, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x52,
+  0x20, 0x3d, 0x20, 0x30, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x69, 0x52, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x0a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e,
+  0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c,
+  0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x3b, 0x20, 0x69, 0x2b, 0x2b, 0x29,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x72, 0x41, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x41,
+  0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77,
+  0x52, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x69, 0x20, 0x2a, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x41, 0x20, 0x3d,
+  0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x69, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x31, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x42, 0x5b, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x69, 0x20, 0x2a, 0x20, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73,
+  0x42, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x70,
+  0x74, 0x72, 0x42, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x69, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31,
+  0x5d, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x2f, 0x2f, 0x28, 0x72, 0x41, 0x20, 0x2b, 0x20, 0x69, 0x41, 0x29, 0x28,
+  0x72, 0x42, 0x20, 0x2b, 0x20, 0x69, 0x42, 0x29, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x3d,
+  0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20,
+  0x3d, 0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72,
+  0x20, 0x3d, 0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6c, 0x61, 0x73, 0x74,
+  0x73, 0x20, 0x3d, 0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b,
+  0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x72, 0x52,
+  0x20, 0x2b, 0x3d, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x2b, 0x20,
+  0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x69, 0x52, 0x20, 0x2b, 0x3d, 0x20, 0x6f, 0x75, 0x74,
+  0x65, 0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74,
+  0x72, 0x52, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72,
+  0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20,
+  0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x52, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x69, 0x52,
+  0x3b, 0x0a, 0x7d, 0x0a, 0x0a, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x5f, 0x6b, 0x6e, 0x6b, 0x0a, 0x28, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x52, 0x2c, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x41, 0x2c, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x42,
+  0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x2c, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f,
+  0x77, 0x73, 0x42, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42,
+  0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x52,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x72,
+  0x6f, 0x77, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73,
+  0x52, 0x20, 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20,
+  0x63, 0x6f, 0x6c, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77,
+  0x52, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62,
+  0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29, 0x3b, 0x20, 0x2f, 0x2f,
+  0x7b, 0x67, 0x70, 0x75, 0x5f, 0x6f, 0x6e, 0x6c, 0x79, 0x7d, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20,
+  0x63, 0x6f, 0x6c, 0x52, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x63, 0x6f,
+  0x6c, 0x52, 0x20, 0x3c, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x3b, 0x20,
+  0x63, 0x6f, 0x6c, 0x52, 0x2b, 0x2b, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x41,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2f, 0x20, 0x72, 0x6f,
+  0x77, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63,
+  0x6f, 0x6c, 0x41, 0x20, 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2f,
+  0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e,
+  0x74, 0x20, 0x72, 0x6f, 0x77, 0x42, 0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77,
+  0x52, 0x20, 0x25, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x42, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x42, 0x20, 0x3d, 0x20,
+  0x63, 0x6f, 0x6c, 0x52, 0x20, 0x25, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42,
+  0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x70, 0x6f, 0x73,
+  0x41, 0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x41, 0x20, 0x2a, 0x20, 0x28,
+  0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b,
+  0x20, 0x28, 0x63, 0x6f, 0x6c, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x70, 0x6f, 0x73, 0x42, 0x20,
+  0x3d, 0x20, 0x72, 0x6f, 0x77, 0x42, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f,
+  0x6c, 0x73, 0x42, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x63, 0x6f, 0x6c, 0x42, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x3b, 0x0a, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x41, 0x20, 0x3d,
+  0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x70, 0x6f, 0x73, 0x41, 0x5d, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x41, 0x20,
+  0x3d, 0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x70, 0x6f, 0x73, 0x41, 0x20,
+  0x2b, 0x20, 0x31, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x42, 0x5b,
+  0x70, 0x6f, 0x73, 0x42, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x42,
+  0x5b, 0x70, 0x6f, 0x73, 0x42, 0x20, 0x2b, 0x20, 0x31, 0x5d, 0x3b, 0x0a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x2f, 0x2f, 0x28,
+  0x72, 0x41, 0x20, 0x2b, 0x20, 0x69, 0x41, 0x29, 0x28, 0x72, 0x42, 0x20,
+  0x2b, 0x20, 0x69, 0x42, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x3d, 0x20, 0x72, 0x41,
+  0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x3d, 0x20, 0x72,
+  0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x20, 0x3d, 0x20,
+  0x69, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66,
+  0x6c, 0x6f, 0x61, 0x74, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x20, 0x3d,
+  0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x72,
+  0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x52,
+  0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69,
+  0x72, 0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72,
+  0x52, 0x5b, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f,
+  0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20,
+  0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x2b,
+  0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7d, 0x0a, 0x7d, 0x0a, 0x0a, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x6b, 0x65, 0x72, 0x6e, 0x65,
+  0x6c, 0x5f, 0x6b, 0x6e, 0x6b, 0x5f, 0x32, 0x78, 0x32, 0x0a, 0x28, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74, 0x72, 0x52,
+  0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62,
+  0x61, 0x6c, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x2a, 0x20, 0x70, 0x74,
+  0x72, 0x41, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x2c,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69,
+  0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x2c, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61,
+  0x74, 0x20, 0x67, 0x61, 0x74, 0x65, 0x30, 0x2c, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x67, 0x61, 0x74, 0x65, 0x31, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20,
+  0x67, 0x61, 0x74, 0x65, 0x32, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67,
+  0x61, 0x74, 0x65, 0x33, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61,
+  0x74, 0x65, 0x34, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61, 0x74,
+  0x65, 0x35, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61, 0x74, 0x65,
+  0x36, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74,
+  0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x67, 0x61, 0x74, 0x65, 0x37,
+  0x0a, 0x29, 0x0a, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x52,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20,
+  0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x3d, 0x20,
+  0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74,
+  0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f,
+  0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29,
+  0x20, 0x2a, 0x20, 0x32, 0x3b, 0x20, 0x2f, 0x2f, 0x7b, 0x67, 0x70, 0x75,
+  0x5f, 0x6f, 0x6e, 0x6c, 0x79, 0x7d, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20,
+  0x3c, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x3b, 0x20, 0x63, 0x6f, 0x6c,
+  0x52, 0x20, 0x2b, 0x3d, 0x20, 0x32, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x72, 0x6f, 0x77, 0x41,
+  0x20, 0x3d, 0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2f, 0x20, 0x32, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e,
+  0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6c, 0x41, 0x20,
+  0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2f, 0x20, 0x32, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x41, 0x20, 0x3d,
+  0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x72, 0x6f, 0x77, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x63, 0x6f, 0x6c, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c,
+  0x6f, 0x61, 0x74, 0x20, 0x69, 0x41, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72,
+  0x41, 0x5b, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f,
+  0x77, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x41,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x31, 0x5d, 0x3b, 0x0a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20,
+  0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x34, 0x3b, 0x20,
+  0x69, 0x2b, 0x2b, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x72, 0x42, 0x2c,
+  0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x73, 0x77, 0x69, 0x74, 0x63, 0x68, 0x20,
+  0x28, 0x69, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61,
+  0x73, 0x65, 0x20, 0x30, 0x3a, 0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x67,
+  0x61, 0x74, 0x65, 0x30, 0x3b, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x67,
+  0x61, 0x74, 0x65, 0x31, 0x3b, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20, 0x31, 0x3a,
+  0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65, 0x32, 0x3b,
+  0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65, 0x33, 0x3b,
+  0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x63, 0x61, 0x73, 0x65, 0x20, 0x32, 0x3a, 0x20, 0x72, 0x42, 0x20, 0x3d,
+  0x20, 0x67, 0x61, 0x74, 0x65, 0x34, 0x3b, 0x20, 0x69, 0x42, 0x20, 0x3d,
+  0x20, 0x67, 0x61, 0x74, 0x65, 0x35, 0x3b, 0x20, 0x62, 0x72, 0x65, 0x61,
+  0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20,
+  0x33, 0x3a, 0x20, 0x72, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65,
+  0x36, 0x3b, 0x20, 0x69, 0x42, 0x20, 0x3d, 0x20, 0x67, 0x61, 0x74, 0x65,
+  0x37, 0x3b, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x2f, 0x2f, 0x28, 0x72, 0x41, 0x20, 0x2b, 0x20, 0x69, 0x41, 0x29,
+  0x28, 0x72, 0x42, 0x20, 0x2b, 0x20, 0x69, 0x42, 0x29, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f,
+  0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x66, 0x69,
+  0x72, 0x73, 0x74, 0x20, 0x3d, 0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x72,
+  0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f,
+  0x61, 0x74, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x3d, 0x20, 0x72,
+  0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73,
+  0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74, 0x20, 0x69, 0x6e, 0x6e, 0x65,
+  0x72, 0x20, 0x3d, 0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x66, 0x6c, 0x6f, 0x61, 0x74,
+  0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x20, 0x3d, 0x20, 0x69, 0x41, 0x20,
+  0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x73, 0x77, 0x69, 0x74, 0x63, 0x68,
+  0x20, 0x28, 0x69, 0x29, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x7b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63,
+  0x61, 0x73, 0x65, 0x20, 0x30, 0x3a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a,
+  0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f,
+  0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a,
+  0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d,
+  0x20, 0x3d, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c,
+  0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20,
+  0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c,
+  0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b,
+  0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20,
+  0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e,
+  0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x62, 0x72, 0x65,
+  0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65,
+  0x20, 0x31, 0x3a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20,
+  0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74,
+  0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70,
+  0x74, 0x72, 0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20,
+  0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29,
+  0x20, 0x2b, 0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73,
+  0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f,
+  0x75, 0x74, 0x65, 0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72,
+  0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20, 0x32, 0x3a,
+  0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72,
+  0x52, 0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72,
+  0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b,
+  0x20, 0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63,
+  0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69, 0x72,
+  0x73, 0x74, 0x20, 0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52,
+  0x5b, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f,
+  0x77, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20,
+  0x28, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f,
+  0x6c, 0x52, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f,
+  0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65,
+  0x5f, 0x74, 0x29, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65,
+  0x72, 0x20, 0x2b, 0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x63, 0x61, 0x73, 0x65, 0x20, 0x33, 0x3a, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52,
+  0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31,
+  0x29, 0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74,
+  0x29, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69,
+  0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28,
+  0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52,
+  0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31,
+  0x29, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x32, 0x29, 0x5d, 0x20, 0x3d, 0x20, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20,
+  0x2b, 0x20, 0x6c, 0x61, 0x73, 0x74, 0x73, 0x3b, 0x0a, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x72, 0x6f, 0x77, 0x52, 0x20,
+  0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x29,
+  0x20, 0x2a, 0x20, 0x28, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a,
+  0x65, 0x5f, 0x74, 0x29, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x28,
+  0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x63, 0x6f, 0x6c, 0x52, 0x20,
+  0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x31, 0x29,
+  0x20, 0x2a, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29, 0x32,
+  0x29, 0x20, 0x2b, 0x20, 0x28, 0x73, 0x69, 0x7a, 0x65, 0x5f, 0x74, 0x29,
+  0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, 0x2b,
+  0x20, 0x69, 0x6e, 0x6e, 0x65, 0x72, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x62, 0x72, 0x65, 0x61, 0x6b, 0x3b, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x0a, 0x20, 0x20,
+  0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x7d, 0x0a, 0x20, 0x20, 0x20, 0x20,
+  0x7d, 0x0a, 0x7d, 0x00
+};
+static unsigned int kernel_gpu_len = 5188;
index 01c0521f30250a6ddc440a39f526f808067a059c..ef9624241171cd045c99e8a9c3bc7bd39493fc8d 100644 (file)
@@ -1,4 +1,4 @@
-#include "qansel.h"
+#include "main.h"
 
 void display_help()
 {
diff --git a/src/main.h b/src/main.h
new file mode 100644 (file)
index 0000000..99fd8ce
--- /dev/null
@@ -0,0 +1,11 @@
+#ifndef __MAIN_H__
+#define __MAIN_H__
+
+#include <stdlib.h>
+#include "context.h"
+#include "openqasm.h"
+
+void display_help();
+void main(int argc, char** argv);
+
+#endif
diff --git a/src/main.o b/src/main.o
new file mode 100644 (file)
index 0000000..9397f00
Binary files /dev/null and b/src/main.o differ
index ee0367ee4dd4a72878790cd5fb8ce5685dacfb96..0873a8352d666210fe1d669374b045d83ea55560 100644 (file)
@@ -1,3 +1,5 @@
+#include "openqasm.h"
+
 int qansel_read_script(char* script, char*** chunksReturn, int** chunkLinesReturn, int* countReturn)
 {
        int scriptSize = strlen(script);
diff --git a/src/openqasm.h b/src/openqasm.h
new file mode 100644 (file)
index 0000000..55f4673
--- /dev/null
@@ -0,0 +1,18 @@
+#ifndef __OPENQASM_H__
+#define __OPENQASM_H__
+
+#include <limits.h>
+#include <stdlib.h>
+#include <regex.h>
+#include "qansel.h"
+#include "context.h"
+#include "bytecode.h"
+
+int qansel_read_script(char* script, char*** chunksReturn, int** chunkLinesReturn, int* countReturn);
+float qansel_parse_float_part(char* neg, char* str);
+int qansel_parse_float(char* str, float* returnFloat);
+int qansel_process_chunk(int index, char* chunk, int line, regmatch_t* regmatches, int* qubitCount, int* bitCount, unsigned char** binary, int* binarySize, QAnselContext* ctx, int samplingshots);
+int qansel_process_chunks(char** chunks, int* associatedLines, int count, unsigned char** retBinary, int* retSize, QAnselContext* ctx, int samplingshots);
+int qanselBuildFromSource(char* source, unsigned char** binary, int* binarySize, QAnselContext* ctx);
+
+#endif
\ No newline at end of file
diff --git a/src/openqasm.o b/src/openqasm.o
new file mode 100644 (file)
index 0000000..46135d2
Binary files /dev/null and b/src/openqasm.o differ
index a1a6782ab664a4a4501a357d4166d2a3123bd34a..eca34d66a5148ab19727ad7ccc961cf6c7b60309 100644 (file)
@@ -2,42 +2,6 @@
 #define __QANSEL_H__
 
 #define QANSEL_MAX_SIZE 1073741824
-#include <stdio.h>
-#include <stdlib.h>
-#include <stdint.h>
-#include <math.h>
-#include <regex.h>
-#include <time.h>
-#include <pthread.h>
-#include <string.h>
-#include <time.h>
-#include <math.h>
-#include <limits.h>
-#if defined(_WIN32) || defined(_WIN64)
-#include <windows.h>
-#elif defined(__linux__)
-#include <unistd.h>
-#include <sys/time.h>
-#elif defined(__APPLE__)
-#include <sys/sysctl.h>
-#endif
-
-#ifdef __SDL2__
-#include <SDL.h>
-#endif
-
-#ifdef __OPENCL__
-#define CL_USE_DEPRECATED_OPENCL_1_2_APIS
-#define CL_TARGET_OPENCL_VERSION 300
-#include <CL/cl.h>
-#define gpuerr(x) if (err != CL_SUCCESS) { fprintf(stderr, "GPU Error on line %i: %s.\n", __LINE__, clGetErrorString(x)); exit(1); }
-cl_platform_id cpx_mtx_platform_id;
-cl_device_id cpx_mtx_device_id;
-cl_context cpx_mtx_context;
-cl_command_queue cpx_mtx_command_queue;
-unsigned char* cpx_mtx_cache = NULL;
-size_t cpx_mtx_cache_len = 0;
-#endif
 
 #define QANSEL_SHOTS_DEFAULT 1000
 #define QANSEL_QUBITS_MAX 16
@@ -94,40 +58,4 @@ size_t cpx_mtx_cache_len = 0;
 #define QANSEL_ALL_CLASSIC 0x3F
 #define QANSEL_ALL 0xFF
 
-typedef struct
-{
-       unsigned char memory_limit;
-       unsigned char hardware_rng;
-       unsigned char verbose;
-       unsigned char optimization_level;
-       unsigned char sampling_bit;
-       float hidden_variable;
-       int display_delay;
-       int sampling_shots;
-       FILE* random_file;
-} QAnselContext;
-
-typedef struct
-{
-       int size;
-       unsigned char bytes[16];
-       int barrier_width;
-       unsigned char barrier[QANSEL_QUBITS_MAX * 2];
-       int next;
-       int prev;
-       int checked;
-       int op;
-       unsigned char ifop[16];
-       int use_ifop;
-} QBytecode;
-
-#include "kernel_cpu.cl"
-#include "kernel_gpu.cl"
-#include "hardware.c"
-#include "complex.c"
-#include "context.c"
-#include "gates.c"
-#include "display.c"
-#include "bytecode.c"
-#include "openqasm.c"
 #endif
\ No newline at end of file