From: miha-q <> Date: Thu, 21 Mar 2024 05:32:40 +0000 (-0400) Subject: Thu Mar 21 01:32:40 AM EDT 2024 X-Git-Url: http://www.foleosoft.com/?a=commitdiff_plain;h=24ae4452eb705d5803df54744b22ea887feaafc4;p=QAnsel.git Thu Mar 21 01:32:40 AM EDT 2024 --- diff --git a/Makefile b/Makefile index 29c84d5..3683099 100644 --- 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`" diff --git a/bin/QAnsel b/bin/QAnsel index 3476638..f6ac158 100755 Binary files a/bin/QAnsel and b/bin/QAnsel differ diff --git a/build.sh b/build.sh index f2c500d..47d8c03 100644 --- 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 index 8218c24..0000000 --- a/examples/belltest_c.txt +++ /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 index 8dab919..0000000 --- a/examples/belltest_q.txt +++ /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 index 0000000..8218c24 --- /dev/null +++ b/examples/chsh_inequality_c.txt @@ -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 index 0000000..815e7c2 --- /dev/null +++ b/examples/chsh_inequality_q.txt @@ -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]; diff --git a/examples/slow.txt b/examples/slow.txt index 65eb3a9..6397512 100644 --- a/examples/slow.txt +++ b/examples/slow.txt @@ -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 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 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 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 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 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 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 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 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 index 0000000..52960f9 --- /dev/null +++ b/src/.kernel.tmp.1 @@ -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 index 0000000..c0f9568 Binary files /dev/null and b/src/.kernel.tmp.2 differ diff --git a/src/bytecode.c b/src/bytecode.c index 0ca140f..6577ebc 100644 --- a/src/bytecode.c +++ b/src/bytecode.c @@ -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 index 0000000..f562277 --- /dev/null +++ b/src/bytecode.h @@ -0,0 +1,145 @@ +#ifndef __BYTECODE_H__ +#define __BYTECODE_H__ + +#include +#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 index 0000000..4379a3f Binary files /dev/null and b/src/bytecode.o differ diff --git a/src/complex.c b/src/complex.c index 6abeeb5..8916788 100644 --- a/src/complex.c +++ b/src/complex.c @@ -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 index 0000000..41efbac --- /dev/null +++ b/src/complex.h @@ -0,0 +1,101 @@ +#ifndef __COMPLEX_H__ +#define __COMPLEX_H__ + +#include +#include +#include +#include +#include +#include "hardware.h" +#include "kernel.h" + +#ifdef __PTHREAD__ +#include +#endif + +#ifdef __OPENCL__ +#define CL_USE_DEPRECATED_OPENCL_1_2_APIS +#define CL_TARGET_OPENCL_VERSION 300 +#include +#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 index 0000000..7183f02 Binary files /dev/null and b/src/complex.o differ diff --git a/src/context.c b/src/context.c index e25abc8..04503bc 100644 --- a/src/context.c +++ b/src/context.c @@ -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 index 0000000..a80670f --- /dev/null +++ b/src/context.h @@ -0,0 +1,26 @@ +#ifndef __CONTEXT_H__ +#define __CONTEXT_H__ + +#include +#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 index 0000000..83f557e Binary files /dev/null and b/src/context.o differ diff --git a/src/display.c b/src/display.c index 489d378..b0519b6 100644 --- a/src/display.c +++ b/src/display.c @@ -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 index 0000000..9c555ad --- /dev/null +++ b/src/display.h @@ -0,0 +1,11 @@ +#ifndef __DISPLAY_H__ +#define __DISPLAY_H__ + +#include +#ifdef __SDL2__ +#include +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 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 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 index 63bbf90..0000000 --- a/src/gates.c +++ /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 diff --git a/src/hardware.c b/src/hardware.c index 130d87e..0283ca2 100644 --- a/src/hardware.c +++ b/src/hardware.c @@ -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 index 0000000..d0cf307 --- /dev/null +++ b/src/hardware.h @@ -0,0 +1,20 @@ +#ifndef __HARDWARE_H__ +#define __HARDWARE_H__ + +#include +#if defined(_WIN32) || defined(_WIN64) +#include +#elif defined(__linux__) +#include +#include +#elif defined(__APPLE__) +#include +#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 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 index de7e2e4..0000000 --- a/src/interface.html +++ /dev/null @@ -1,40 +0,0 @@ - - - - - - - - -
HXYZ
- - - - - - - - - - - - - - - - - - - - - - - - - - - - -
Q0---
Q1---
Q2
Q3---
Q4---
- - \ No newline at end of file diff --git a/src/kernel.cl b/src/kernel.cl index dcfdf50..a5731c3 100644 --- a/src/kernel.cl +++ b/src/kernel.cl @@ -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 index 0000000..f71b799 --- /dev/null +++ b/src/kernel.h @@ -0,0 +1,48 @@ +#ifndef __KERNEL_H__ +#define __KERNEL_H__ + +#include + +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 index 0000000..842f3c5 --- /dev/null +++ b/src/kernel_cpu.c @@ -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 index 0000000..842f3c5 --- /dev/null +++ b/src/kernel_cpu.cl @@ -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 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 index 0000000..c489f41 --- /dev/null +++ b/src/kernel_gpu.c @@ -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 index 0000000..c489f41 --- /dev/null +++ b/src/kernel_gpu.cl @@ -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/main.c b/src/main.c index 01c0521..ef96242 100644 --- a/src/main.c +++ b/src/main.c @@ -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 index 0000000..99fd8ce --- /dev/null +++ b/src/main.h @@ -0,0 +1,11 @@ +#ifndef __MAIN_H__ +#define __MAIN_H__ + +#include +#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 index 0000000..9397f00 Binary files /dev/null and b/src/main.o differ diff --git a/src/openqasm.c b/src/openqasm.c index ee0367e..0873a83 100644 --- a/src/openqasm.c +++ b/src/openqasm.c @@ -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 index 0000000..55f4673 --- /dev/null +++ b/src/openqasm.h @@ -0,0 +1,18 @@ +#ifndef __OPENQASM_H__ +#define __OPENQASM_H__ + +#include +#include +#include +#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 index 0000000..46135d2 Binary files /dev/null and b/src/openqasm.o differ diff --git a/src/qansel.h b/src/qansel.h index a1a6782..eca34d6 100644 --- a/src/qansel.h +++ b/src/qansel.h @@ -2,42 +2,6 @@ #define __QANSEL_H__ #define QANSEL_MAX_SIZE 1073741824 -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#if defined(_WIN32) || defined(_WIN64) -#include -#elif defined(__linux__) -#include -#include -#elif defined(__APPLE__) -#include -#endif - -#ifdef __SDL2__ -#include -#endif - -#ifdef __OPENCL__ -#define CL_USE_DEPRECATED_OPENCL_1_2_APIS -#define CL_TARGET_OPENCL_VERSION 300 -#include -#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