From 1724ae38eaea96714d79e59a25ed704dedc00cda Mon Sep 17 00:00:00 2001 From: miha-q <> Date: Sun, 10 Mar 2024 17:55:40 -0400 Subject: [PATCH] Sun Mar 10 05:55:40 PM EDT 2024 --- examples/slow.txt | 24 +- examples/slower.txt | 34 +++ examples/slowest.txt | 24 ++ src/.kernel.tmp.1 | 37 ++- src/.kernel.tmp.2 | Bin 4428 -> 5187 bytes src/QAnsel.c | 8 +- src/bytecode.c | 316 +++++++++++---------- src/kernel.cl | 37 ++- src/kernel_cpu.cl | 37 ++- src/kernel_gpu.cl | 648 ++++++++++++++++++++++++------------------- src/oldex.mm | 127 --------- src/openqasm.c | 16 +- src/qansel.h | 15 +- 13 files changed, 678 insertions(+), 645 deletions(-) create mode 100644 examples/slower.txt create mode 100644 examples/slowest.txt delete mode 100644 src/oldex.mm diff --git a/examples/slow.txt b/examples/slow.txt index b01c9ad..8767db3 100644 --- a/examples/slow.txt +++ b/examples/slow.txt @@ -1,6 +1,22 @@ -//designed to be slow -qreg q[14]; -x q[0]; +//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[13]; +h q[0]; x q[1]; x q[2]; x q[3]; @@ -13,6 +29,4 @@ x q[9]; x q[10]; x q[11]; x q[12]; -x q[13]; - born q; \ No newline at end of file diff --git a/examples/slower.txt b/examples/slower.txt new file mode 100644 index 0000000..bab3c85 --- /dev/null +++ b/examples/slower.txt @@ -0,0 +1,34 @@ +//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. If +// queueing is enabled, this program will run +// almost as fast as if there is just a single +// instruction. + +qreg q[14]; +h q[0]; +x q[1]; +x q[2]; +x q[3]; +x q[4]; +x q[5]; +x q[6]; +x q[7]; +x q[8]; +x q[9]; +x q[10]; +x q[11]; +x q[12]; +x q[13]; +born q; \ No newline at end of file diff --git a/examples/slowest.txt b/examples/slowest.txt new file mode 100644 index 0000000..25f03b0 --- /dev/null +++ b/examples/slowest.txt @@ -0,0 +1,24 @@ +//This will basically never finish unless you run it +// with queueing and threading enabled. It also will +// likely crash if GPU is enabled due to running +// out of GPU memory. It uses so much RAM you will +// likely need at least 32 gigabytes of swap space +// for it not to crash. +qreg q[16]; +h q[0]; +x q[1]; +x q[2]; +x q[3]; +x q[4]; +x q[5]; +x q[6]; +x q[7]; +x q[8]; +x q[9]; +x q[10]; +x q[11]; +x q[12]; +x q[13]; +x q[14]; +x q[15]; +born q; \ No newline at end of file diff --git a/src/.kernel.tmp.1 b/src/.kernel.tmp.1 index 1da9581..2a8664f 100644 --- a/src/.kernel.tmp.1 +++ b/src/.kernel.tmp.1 @@ -16,17 +16,13 @@ __kernel void kernel_dot float rR = 0; float iR = 0; - const int posR = rowR * (colsR * 2) + (colR * 2); for (int i = 0; i < colsA; i++) { - const int posA = rowR * (colsA * 2) + (i * 2); - const int posB = i * (colsB * 2) + (colR * 2); - - const float rA = ptrA[posA]; - const float iA = ptrA[posA + 1]; - const float rB = ptrB[posB]; - const float iB = ptrB[posB + 1]; + 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; @@ -37,8 +33,8 @@ __kernel void kernel_dot rR += first + lasts; iR += outer + inner; } - ptrR[rowR * (colsR * 2) + (colR * 2)] = rR; - ptrR[rowR * (colsR * 2) + (colR * 2) + 1] = iR; + 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 @@ -104,9 +100,8 @@ __kernel void kernel_knk_2x2 { const int rowA = rowR / 2; const int colA = colR / 2; - const int posA = rowA * (colsA * 2) + (colA * 2); - const float rA = ptrA[posA]; - const float iA = ptrA[posA + 1]; + 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++) { @@ -127,20 +122,20 @@ __kernel void kernel_knk_2x2 switch (i) { case 0: - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts; - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner; + 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[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts; - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner; + 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[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts; - ptrR[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner; + 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[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts; - ptrR[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner; + 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 index 1e791cc72a16245bb1931b4a70dfb487507ea110..fd607072464755b25e62b8061b835bfd592badce 100644 GIT binary patch literal 5187 zcmds4OK+P%5bl{@v6n~?N3y|rNO5z3pWtg1A;%%k$_OGLP20%-y)&~97T6_;Ory43 zZ05D|{B}sweO47&&c2uYo{9hIec5=v$KXlRU0&X8a`upyo92`~HP!me#fhDqL~yoc zQ8$bi4XetRT8g*CZJ7^=C^=9^$;1o1S8}DbfWm8bBh0UQOjJT9DfgM6iM(Cb$cR&Z z?=ONCo6lc&PtR#tO%0Yv2rX%kp;u=&Ot>>+B-c6vZ3|4{ z2E>9V38okmL<(IqTPJoaGhQ}~b1ke-X}41;%mAEm@VwN9iTh&zA7CQ30RpPXWec1* zM+id-B51DqLhr}O2ZdrZfrxi#c`dPQE#6#kChrGC8ai=qy4=nRih)rSBinOwelQ5q z8csM_I~lcxVr1<^S<}PNChMuJrE@q|KFGuSDHCu?F5!#FUe{mHl#V1y6aQ#U@nTG% z9pKRc7nEs$CNF2~ZD2*GT-G!&lIOyj&d5IlXjlY%2s@|?WZ4jdbn?Z^Ep&hb@!q2j zkP@N^D=v;;OB(%%j=-0;IU&5^5=AX)?qFhN*ln6@VJC-na%3kj?c{Si`NB@Vw3Dyi z+Vil6eW=2KO*|V2P!|mMJ+jr^`KaCB1!#C5M)nnLh-uOzY|FNI1RfgC1|lPi9wCvr z6%oO@=t+c9!cpKFehno`6q0t#Ii2g)06CppW9L0L_Fc@;t41AslZi;qV25gP7CH3` zZ?=!HLbt0GlDq0i*>38LEv{j|f@!5efD_gKaK-MbY_sp$)@q@l78z>Gaka=$!}!!` z7vpM6L+#Q~yRxbQ#-N|-8O3N6GH)2?;@KB;E}wlt=K|Un94w(Ipq`HwwUYs5L@bU* zE5rV(KR*@$On&D>$+Dg*T^OLDt{%0K3q-kszn|rf5v1#4GTj<)`_D~mNJaT;65DfU z!o)_`A5X0C#%M{n?*^at2Kx_p?1b)2nAqhfC$`rjQ$WEniS4;_d}58uMSSHSyu+8@ D#Kf67 delta 539 zcmX@CaYkuFH>0Qm5G3d46_+Sv=9MTEi zAO$T2jpY2CVn-muNK-*uK_e5)ntYK>WO5!myEupoR^bF!0aPEfxt~3Zak2=D*W^ze z$}B}e)?AZca%fBr;N;lM%esIKXt>qn|4eL?SF>AfuI4Fa(kRL=4+0u*pb0Y%WHrP_ zknwQ2$pQTOszfQC?9Hh#nO{VHvJNlzWIX|CVyrL(Sx>ALhMI6IY6R6_#)FLjhX@fq bz^@o?ybdqt QANSEL_MODE_METAL_THREADED) + if (optimizationSettings < QANSEL_MODE_BARE || optimizationSettings > (QANSEL_MODE_THREADED | QANSEL_MODE_METAL | QANSEL_MODE_QUEUE)) { fprintf(stderr, "QAnsel: Invalid optimization settings.\n"); exit(1); @@ -54,7 +54,7 @@ void main(int argc, char** argv) } } - if (QANSEL_MODE == QANSEL_MODE_METAL || QANSEL_MODE == QANSEL_MODE_METAL_THREADED) + if (QANSEL_MODE & QANSEL_MODE_METAL) { if (!cpx_mtx_begin()) { @@ -82,7 +82,7 @@ void main(int argc, char** argv) free(bytecode); if (QANSEL_RANDOM_FILE != NULL) fclose(QANSEL_RANDOM_FILE); - if (QANSEL_MODE == QANSEL_MODE_METAL || QANSEL_MODE == QANSEL_MODE_METAL_THREADED) cpx_mtx_clean(); + if (QANSEL_MODE & QANSEL_MODE_METAL) cpx_mtx_clean(); return; } \ No newline at end of file diff --git a/src/bytecode.c b/src/bytecode.c index c593ca2..d4e2c3c 100644 --- a/src/bytecode.c +++ b/src/bytecode.c @@ -197,25 +197,25 @@ float* qansel_unitary(float theta, float phi, float lambda) void qansel_queue_init ( - cpx_mtx_t* queueVectors, + cpx_mtx_t* queueVector, int qubitCount, int dofree ) { for (int i = 0; i < qubitCount; i++) { - if (dofree) free(queueVectors[i].ptr); - int vectLen = queueVectors[i].rows * (queueVectors[i].cols * 2) + queueVectors[i].cols * 2; - queueVectors[i].rows = 2; - queueVectors[i].cols = 2; - queueVectors[i].ptr = malloc(vectLen); - memcpy(queueVectors[i].ptr, Identity, vectLen); + if (dofree) free(queueVector[i].ptr); + int vectLen = queueVector[i].rows * (queueVector[i].cols * 2) + queueVector[i].cols * 2; + queueVector[i].rows = 2; + queueVector[i].cols = 2; + queueVector[i].ptr = malloc(vectLen); + memcpy(queueVector[i].ptr, Identity, vectLen); } } void qansel_instruction_queue ( - cpx_mtx_t** queueVectors, + cpx_mtx_t** queueVector, int qubitCount, unsigned char instr, unsigned char index, @@ -252,14 +252,14 @@ void qansel_instruction_queue gate.rows = 2; gate.cols = 2; gate.ptr = gate_ptr; - tmp.rows = 1; + tmp.rows = 2; tmp.cols = 2; - tmp.ptr = malloc(tmp.rows * (tmp.cols * 2) * sizeof(float)); - cpx_mtx_dot(tmp.ptr, queueVectors[index]->ptr, gate_ptr, queueVectors[index]->rows, queueVectors[index]->cols, 2, 2); - free(queueVectors[index]->ptr); - queueVectors[index]->rows = tmp.rows; - queueVectors[index]->cols = tmp.cols; - queueVectors[index]->ptr = tmp.ptr; + tmp.ptr = malloc((tmp.rows * 2) * (tmp.cols * 2) * sizeof(float)); + cpx_mtx_dot(tmp.ptr, (*queueVector)[qubit].ptr, gate_ptr, (*queueVector)[qubit].rows, (*queueVector)[qubit].cols, 2, 2); + free((*queueVector)[qubit].ptr); + (*queueVector)[qubit].rows = tmp.rows; + (*queueVector)[qubit].cols = tmp.cols; + (*queueVector)[qubit].ptr = tmp.ptr; if (needToFreeGate) free(gate_ptr); } @@ -272,12 +272,10 @@ void qansel_instruction float arg0, float arg1, float arg2, - //if using gate queueing - cpx_mtx_t* queueVectors, - int queueStatus + //pass in this to flush the queue + cpx_mtx_t** queueVector ) { - if (queueStatus == 0) queueVectors = NULL; cpx_mtx_t tmp; cpx_mtx_t gate; gate.rows = 2; @@ -307,23 +305,23 @@ void qansel_instruction cpx_t n; cpx_mtx_t filter; cpx_mtx_init(&filter, 2, 2); - unsigned char qubit = qubitCount - (queueVectors == NULL ? index : 0) - 1; - if (qubit == 0 && queueVectors == NULL) + unsigned char qubit = qubitCount - (queueVector == NULL ? index : 0) - 1; + if (qubit == 0 && queueVector == NULL) { memcpy(filter.ptr, gate_ptr, 8 * sizeof(float)); } - else if (queueVectors == NULL) + else if (queueVector == NULL) { memcpy(filter.ptr, Identity, 8 * sizeof(float)); } else { - memcpy(filter.ptr, queueVectors[0].ptr, 8 * sizeof(float)); + memcpy(filter.ptr, (*queueVector)[0].ptr, 8 * sizeof(float)); } for (unsigned char i = 1; i < qubitCount; i++) { - if (index != 0x0F && queueVectors == NULL) + if (index != QANSEL_ALL_QUANTUM && queueVector == NULL) { if (qubit == i) { @@ -334,19 +332,23 @@ void qansel_instruction gate.ptr = Identity; } } - else if (queueVectors == NULL) + else if (queueVector == NULL) { gate.ptr = gate_ptr; } else { - gate.ptr = queueVectors[i].ptr; + gate.ptr = (*queueVector)[i].ptr; } tmp.rows = filter.rows * gate.rows; tmp.cols = filter.cols * gate.cols; - tmp.ptr = malloc(tmp.rows * (tmp.cols * 2) * sizeof(float)); - + tmp.ptr = malloc((size_t)tmp.rows * ((size_t)tmp.cols * (size_t)2) * (size_t)sizeof(float)); + if (tmp.ptr == NULL) + { + fprintf(stderr, "QAnsel: Ran out of memory.\n"); + exit(1); + } #ifdef SPEED_TEST printf("(%ix%i);(%ix%i) (knk)\n", tmp.rows, tmp.cols, gate.rows, gate.cols); unsigned long int us1, us2; @@ -380,11 +382,11 @@ void qansel_instruction //us2 = get_time(); //printf("\tTranspose: %lu\n", us2 - us1); #else - if (QANSEL_MODE == QANSEL_MODE_METAL && tmp.cols >= 64) + if ((QANSEL_MODE & QANSEL_MODE_METAL) && !(QANSEL_MODE & QANSEL_MODE_THREADED) && tmp.cols >= 64) { cpx_mtx_knk_metal_2x2(tmp.ptr, filter.ptr, gate.ptr, filter.rows, filter.cols, gate.rows, gate.cols); } - else if ((QANSEL_MODE == QANSEL_MODE_THREADED || QANSEL_MODE == QANSEL_MODE_METAL_THREADED) && tmp.cols >= 64) + else if ((QANSEL_MODE & QANSEL_MODE_THREADED) && tmp.cols >= 64) { cpx_mtx_knk_threads_2x2(tmp.ptr, filter.ptr, gate.ptr, filter.rows, filter.cols, gate.rows, gate.cols); } @@ -418,11 +420,11 @@ void qansel_instruction us2 = get_time(); printf("\tBare: %lu\n", us2 - us1); #else - if ((QANSEL_MODE == QANSEL_MODE_METAL || QANSEL_MODE == QANSEL_MODE_METAL_THREADED) && tmp.cols >= 64) + if ((QANSEL_MODE & QANSEL_MODE_METAL) && tmp.cols >= 64) { cpx_mtx_dot_metal(tmp.ptr, stateVector->ptr, filter.ptr, stateVector->rows, stateVector->cols, filter.rows, filter.cols); } - else if (QANSEL_MODE == QANSEL_MODE_THREADED && tmp.cols >= 64) + else if ((QANSEL_MODE & QANSEL_MODE_THREADED) && tmp.cols >= 64) { cpx_mtx_dot_threads(tmp.ptr, stateVector->ptr, filter.ptr, stateVector->rows, stateVector->cols, filter.rows, filter.cols); } @@ -507,8 +509,8 @@ int qansel_get_instruction_bitmax(unsigned char* ptr, int offset, int* bitmax, i case QANSEL_INSTRUCTION_DENSITY: case QANSEL_INSTRUCTION_BORN: a0 = ptr[offset + 1]; - if (a0 > 0x0D && a0 != 0x0F) return 0; - if (a0 != 0x0F) *qbitmax = a0 + 1; + if (a0 > QANSEL_QBOUND_UPPER && a0 != QANSEL_ALL_QUANTUM) return 0; + if (a0 != QANSEL_ALL_QUANTUM) *qbitmax = a0 + 1; return 1; case QANSEL_INSTRUCTION_SAMPLE: case QANSEL_INSTRUCTION_IF_E: @@ -518,23 +520,23 @@ int qansel_get_instruction_bitmax(unsigned char* ptr, int offset, int* bitmax, i case QANSEL_INSTRUCTION_IF_L: case QANSEL_INSTRUCTION_IF_LE: a0 = ptr[offset + 1]; - if ((a0 > 0x1D || a0 < 0x10) && a0 != 0x1F) return 0; - if (a0 != 0x1F) *bitmax = (a0 - 0x10) + 1; + if ((a0 > QANSEL_CBOUND_UPPER || a0 < QANSEL_CBOUND_LOWER) && a0 != QANSEL_ALL_CLASSIC) return 0; + if (a0 != QANSEL_ALL_CLASSIC) *bitmax = (a0 - QANSEL_CBOUND_LOWER) + 1; return 1; case QANSEL_INSTRUCTION_RESET: case QANSEL_INSTRUCTION_PRINT: a0 = ptr[offset + 1]; - if (a0 == 0xFF) return 1; - if (a0 < 0x10) + if (a0 == QANSEL_ALL) return 1; + if (a0 < QANSEL_CBOUND_LOWER) { - if (a0 > 0x0D && a0 != 0x0F) return 0; - if (a0 != 0x0F) *qbitmax = a0 + 1; + if (a0 > QANSEL_QBOUND_UPPER && a0 != QANSEL_ALL_QUANTUM) return 0; + if (a0 != QANSEL_ALL_QUANTUM) *qbitmax = a0 + 1; return 1; } else { - if ((a0 > 0x1D || a0 < 0x10) && a0 != 0x1F) return 0; - if (a0 != 0x1F) *bitmax = (a0 - 0x10) + 1; + if ((a0 > QANSEL_CBOUND_UPPER || a0 < QANSEL_CBOUND_LOWER) && a0 != QANSEL_ALL_CLASSIC) return 0; + if (a0 != QANSEL_ALL_CLASSIC) *bitmax = (a0 - QANSEL_CBOUND_LOWER) + 1; return 1; } return 0; @@ -542,8 +544,8 @@ int qansel_get_instruction_bitmax(unsigned char* ptr, int offset, int* bitmax, i case QANSEL_INSTRUCTION_SWAP: a0 = ptr[offset + 1]; a1 = ptr[offset + 2]; - if (a0 > 0x0D) return 0; - if (a1 > 0x0D) return 0; + if (a0 > QANSEL_QBOUND_UPPER) return 0; + if (a1 > QANSEL_QBOUND_UPPER) return 0; *qbitmax = (a0 > a1 ? a0 : a1) + 1; return 1; case QANSEL_INSTRUCTION_CCX: @@ -551,16 +553,16 @@ int qansel_get_instruction_bitmax(unsigned char* ptr, int offset, int* bitmax, i a0 = ptr[offset + 1]; a1 = ptr[offset + 2]; a2 = ptr[offset + 3]; - if (a0 > 0x0D || a1 > 0x0D || a2 > 0x0D) return 0; + if (a0 > QANSEL_QBOUND_UPPER || a1 > QANSEL_QBOUND_UPPER || a2 > QANSEL_QBOUND_UPPER) return 0; *qbitmax = ((a0 > a1) && (a0 > a2) ? a0 : ((a1 > a0) && (a1 > a2) ? a1 : a2)) + 1; return 1; case QANSEL_INSTRUCTION_MEASURE: a0 = ptr[offset + 1]; a1 = ptr[offset + 2]; - if (a0 > 0x0D) return 0; - if (a1 > 0x1D || a1 < 0x10) return 0; + if (a0 > QANSEL_QBOUND_UPPER) return 0; + if (a1 > QANSEL_CBOUND_UPPER || a1 < QANSEL_CBOUND_LOWER) return 0; *qbitmax = a0 + 1; - *bitmax = (a1 - 0x10) + 1; + *bitmax = (a1 - QANSEL_CBOUND_LOWER) + 1; return 1; case QANSEL_INSTRUCTION_RAND: case QANSEL_INSTRUCTION_HVAR: @@ -613,7 +615,7 @@ int qansel_get_instruction_size(unsigned char instr) void qansel_born(cpx_mtx_t* stateVector, int PC, int qubitCount, unsigned char q0) { unsigned int qubitCountPow2 = (unsigned int)pow(2, qubitCount); - if (q0 == 0x0F) + if (q0 == QANSEL_ALL_QUANTUM) { for (unsigned int j = 0; j < qubitCountPow2; j++) { @@ -628,7 +630,7 @@ void qansel_born(cpx_mtx_t* stateVector, int PC, int qubitCount, unsigned char q printf(": %.1f%%\n", cpx_magsqr(&n) * 100); } } - else if (q0 <= 0x0D) + else if (q0 <= QANSEL_QBOUND_UPPER) { float prob = 0; for (unsigned int j = 0; j < qubitCountPow2; j++) @@ -648,13 +650,13 @@ void qansel_born(cpx_mtx_t* stateVector, int PC, int qubitCount, unsigned char q void qansel_density_or_print(cpx_mtx_t* stateVector, unsigned char* bitVector, unsigned char density, int bitCount, int qubitCount, unsigned char a0) { unsigned int qubitCountPow2 = (unsigned int)pow(2, qubitCount); - if (a0 == 0x0F || a0 == 0x1F || a0 == 0xFF) + if (a0 == QANSEL_ALL_QUANTUM || a0 == QANSEL_ALL_CLASSIC || a0 == QANSEL_ALL) { - if (a0 == 0x0F || a0 == 0xFF) + if (a0 == QANSEL_ALL_QUANTUM || a0 == QANSEL_ALL) { printf("[ "); cpx_mtx_print(stateVector); printf(" ]\n"); } - if (a0 == 0x1F || a0 == 0xFF) + if (a0 == QANSEL_ALL_CLASSIC || a0 == QANSEL_ALL) { for (int32_t j = bitCount - 1; j >= 0; j--) { @@ -663,9 +665,9 @@ void qansel_density_or_print(cpx_mtx_t* stateVector, unsigned char* bitVector, u putchar('\n'); } } - else if (a0 >= 0x10 && a0 <= 0x1D) + else if (a0 >= QANSEL_CBOUND_LOWER && a0 <= QANSEL_CBOUND_UPPER) { - putchar('0' + bitVector[a0 - 0x10]); + putchar('0' + bitVector[a0 - QANSEL_CBOUND_LOWER]); putchar('\n'); } else @@ -764,7 +766,7 @@ int qansel_get_int(unsigned char* program, int offset) void qansel_reset(cpx_mtx_t* stateVector, unsigned char* bitVector, int qubitCount, int bitCount, unsigned char q0) { unsigned int qubitCountPow2 = (unsigned int)pow(2, qubitCount); - if (q0 == 0xFF) + if (q0 == QANSEL_ALL) { cpx_mtx_set2(stateVector, 0, 0, 1, 0); for (unsigned int j = 1; j < qubitCountPow2; j++) @@ -776,7 +778,7 @@ void qansel_reset(cpx_mtx_t* stateVector, unsigned char* bitVector, int qubitCou bitVector[j] = 0; } } - else if (q0 == 0x0F) + else if (q0 == QANSEL_ALL_QUANTUM) { cpx_mtx_set2(stateVector, 0, 0, 1, 0); for (unsigned int j = 1; j < qubitCountPow2; j++) @@ -784,24 +786,24 @@ void qansel_reset(cpx_mtx_t* stateVector, unsigned char* bitVector, int qubitCou cpx_mtx_set2(stateVector, 0, j, 0, 0); } } - else if (q0 == 0x1F) + else if (q0 == QANSEL_ALL_CLASSIC) { for (unsigned char j = 0; j < bitCount; j++) { bitVector[j] = 0; } } - else if (q0 <= 0x0D) + else if (q0 <= QANSEL_QBOUND_UPPER) { unsigned char bit = qansel_measure(stateVector, qubitCount, q0); if (bit) { - qansel_instruction(stateVector, qubitCount, QANSEL_INSTRUCTION_X, q0, 0, 0, 0, NULL, 0); + qansel_instruction(stateVector, qubitCount, QANSEL_INSTRUCTION_X, q0, 0, 0, 0, NULL); } } - else if (q0 >= 0x10 && q0 <= 0x1D) + else if (q0 >= QANSEL_CBOUND_LOWER && q0 <= QANSEL_CBOUND_UPPER) { - bitVector[q0 - 0x10] = 0; + bitVector[q0 - QANSEL_CBOUND_LOWER] = 0; } } @@ -809,7 +811,7 @@ unsigned char qansel_compare(unsigned char* bitVector, int bitCount, int PC, uns { unsigned char ret = 0; short val; - if (a0 == 0x1F) + if (a0 == QANSEL_ALL_CLASSIC) { val = 0; for (int32_t j = bitCount - 1; j >= 0; j--) @@ -820,9 +822,9 @@ unsigned char qansel_compare(unsigned char* bitVector, int bitCount, int PC, uns if (val > op) ret |= QANSEL_FLAGS_GREATER; if (val < op) ret |= QANSEL_FLAGS_LESSER; } - else if (a0 >= 0x10 && a0 <= 0x1D) + else if (a0 >= QANSEL_CBOUND_LOWER && a0 <= QANSEL_CBOUND_UPPER) { - val = bitVector[a0 - 0x10]; + val = bitVector[a0 - QANSEL_CBOUND_LOWER]; if (val == op) ret |= QANSEL_FLAGS_EQUAL; if (val > op) ret |= QANSEL_FLAGS_GREATER; if (val < op) ret |= QANSEL_FLAGS_LESSER; @@ -836,19 +838,19 @@ int qansel_crawl(unsigned char* program, int programSize, int* qubitCount, int* int PC = 0; *qubitCount = 0; *bitCount = 0; - *sample = 0xFF; + *sample = QANSEL_ALL; while (PC < programSize) { int next = qansel_get_instruction_size(program[PC]); if (program[PC] == QANSEL_INSTRUCTION_SAMPLE) { - if ((program[PC + 1] < 0x10 || program[PC + 1] > 0x1D) && program[PC + 1] != 0x1F) + if ((program[PC + 1] < QANSEL_CBOUND_LOWER || program[PC + 1] > QANSEL_CBOUND_UPPER) && program[PC + 1] != QANSEL_ALL_CLASSIC) { if (QANSEL_VERBOSE) fprintf(stderr, "QAnsel (%04X): Invalid index.\n", PC); } else { - *sample = program[PC + 1] - 0x10; + *sample = program[PC + 1] - QANSEL_CBOUND_LOWER; } } if (next == 0) @@ -874,14 +876,13 @@ int qansel_crawl(unsigned char* program, int programSize, int* qubitCount, int* void qansel_run(unsigned char* program, int programSize, int qubitCount, int bitCount, unsigned char* outputBitVector) { - int useQueue = 0; + int useQueue = (QANSEL_MODE & QANSEL_MODE_QUEUE) ? 1 : 0; int PC = 0; unsigned int qubitCountPow2 = (unsigned int)pow(2, qubitCount); unsigned char bitVector[bitCount]; memset(bitVector, 0, bitCount); cpx_mtx_t stateVector; cpx_mtx_t* queueVector = NULL; - int queueStatus = 0; cpx_mtx_init(&stateVector, 1, qubitCountPow2); cpx_mtx_set2(&stateVector, 0, 0, 1, 0); if (useQueue) @@ -889,14 +890,18 @@ void qansel_run(unsigned char* program, int programSize, int qubitCount, int bit queueVector = malloc(qubitCount * sizeof(cpx_mtx_t)); for (int i = 0; i < qubitCount; i++) { - cpx_mtx_init(&(queueVector[i]), 1, 2); - cpx_mtx_set2(&(queueVector[i]), 0, 0, 1, 0); + queueVector[i].ptr = malloc(sizeof(Identity)); + memcpy(queueVector[i].ptr, Identity, sizeof(Identity)); + queueVector[i].rows = 2; + queueVector[i].cols = 2; } } if (QANSEL_USE_DISPLAY) { QANSEL_USE_DISPLAY = display(&stateVector, qubitCount, QANSEL_USE_DISPLAY); } unsigned char skip = 0, a0 = 0, a1 = 0, a2 = 0; unsigned char flags = 0; unsigned short tmp = 0; + unsigned char queueFlushed = 1; + float f0, f1, f2; while (PC < programSize) { @@ -910,9 +915,40 @@ void qansel_run(unsigned char* program, int programSize, int qubitCount, int bit { unsigned char instr = program[PC]; //printf("-----------------------------------\n"); - //qansel_density_or_print(&stateVector, bitVector, 0, bitCount, qubitCount, 0x0F); - //qansel_density_or_print(&stateVector, bitVector, 0, bitCount, qubitCount, 0x1F); + //qansel_density_or_print(&stateVector, bitVector, 0, bitCount, qubitCount, QANSEL_ALL_QUANTUM); + //qansel_density_or_print(&stateVector, bitVector, 0, bitCount, qubitCount, QANSEL_ALL_CLASSIC); //printf("%s(%i, %i, %i)(%f, %f, %f)\n", qansel_instruction_to_string(instr), program[PC+1], program[PC+2], program[PC], qansel_get_float(program, PC + 2), qansel_get_float(program, PC + 2 + sizeof(float)), qansel_get_float(program, PC + 2) + sizeof(float) * 2); + + //flush the queue if any non-single qubit instructions are called + if (useQueue) + { + switch (instr) + { + case QANSEL_INSTRUCTION_X: + case QANSEL_INSTRUCTION_Y: + case QANSEL_INSTRUCTION_Z: + case QANSEL_INSTRUCTION_H: + case QANSEL_INSTRUCTION_S: + case QANSEL_INSTRUCTION_T: + case QANSEL_INSTRUCTION_RX: + case QANSEL_INSTRUCTION_RY: + case QANSEL_INSTRUCTION_RZ: + case QANSEL_INSTRUCTION_U1: + case QANSEL_INSTRUCTION_U2: + case QANSEL_INSTRUCTION_U3: + queueFlushed = 0; + break; + default: + qansel_instruction(&stateVector, qubitCount, instr, 0, 0, 0, 0, &queueVector); + for (int i = 0; i < qubitCount; i++) + { + memcpy(queueVector[i].ptr, Identity, sizeof(Identity)); + } + queueFlushed = 1; + break; + } + } + switch (instr) { case QANSEL_INSTRUCTION_X: @@ -921,150 +957,142 @@ void qansel_run(unsigned char* program, int programSize, int qubitCount, int bit case QANSEL_INSTRUCTION_H: case QANSEL_INSTRUCTION_S: case QANSEL_INSTRUCTION_T: - a0 = program[PC + 1]; - qansel_instruction(&stateVector, qubitCount, instr, program[PC + 1], 0, 0, 0, queueVector, queueStatus); - break; case QANSEL_INSTRUCTION_RX: - a0 = program[PC + 1]; - qansel_instruction - ( - &stateVector, qubitCount, instr, a0, - M_PI / 2, -M_PI / 2, qansel_get_float(program, PC + 2) - (M_PI / 2), - queueVector, queueStatus - ); - break; case QANSEL_INSTRUCTION_RY: - case QANSEL_INSTRUCTION_U1: - a0 = program[PC + 1]; - qansel_instruction - ( - &stateVector, qubitCount, instr, a0, - qansel_get_float(program, PC + 2), 0, 0, - queueVector, queueStatus - ); - break; case QANSEL_INSTRUCTION_RZ: - a0 = program[PC + 1]; - qansel_instruction - ( - &stateVector, qubitCount, instr, a0, - 0, 0, qansel_get_float(program, PC + 2), - queueVector, queueStatus - ); - break; + case QANSEL_INSTRUCTION_U1: case QANSEL_INSTRUCTION_U2: - a0 = program[PC + 1]; - qansel_instruction - ( - &stateVector, qubitCount, instr, a0, - qansel_get_float(program, PC + 2), - qansel_get_float(program, PC + 2 + sizeof(float)), - 0, queueVector, queueStatus - ); - break; case QANSEL_INSTRUCTION_U3: + switch (instr) + { + case QANSEL_INSTRUCTION_RX: + f0 = M_PI / 2; + f1 = -M_PI / 2; + f2 = qansel_get_float(program, PC + 2) - (M_PI / 2); + break; + case QANSEL_INSTRUCTION_RY: + case QANSEL_INSTRUCTION_U1: + f0 = qansel_get_float(program, PC + 2); + f1 = 0; + f2 = 0; + break; + case QANSEL_INSTRUCTION_RZ: + f0 = 0; + f1 = 0; + f2 = qansel_get_float(program, PC + 2); + break; + case QANSEL_INSTRUCTION_U2: + f0 = qansel_get_float(program, PC + 2); + f1 = qansel_get_float(program, PC + 2 + sizeof(float)); + f2 = 0; + break; + case QANSEL_INSTRUCTION_U3: + f0 = qansel_get_float(program, PC + 2); + f1 = qansel_get_float(program, PC + 2 + sizeof(float)); + f2 = qansel_get_float(program, PC + 2 + sizeof(float) * 2); + break; + default: + f0 = 0; + f1 = 0; + f2 = 0; + break; + } a0 = program[PC + 1]; - qansel_instruction - ( - &stateVector, qubitCount, instr, a0, - qansel_get_float(program, PC + 2), - qansel_get_float(program, PC + 2 + sizeof(float)), - qansel_get_float(program, PC + 2 + sizeof(float) * 2), - queueVector, queueStatus - ); - break; + if (useQueue) qansel_instruction_queue(&queueVector, qubitCount, instr, a0, f0, f1, f2); + else qansel_instruction(&stateVector, qubitCount, instr, a0, f0, f1, f2, NULL); + break; case QANSEL_INSTRUCTION_CX: a0 = program[PC + 1]; a1 = program[PC + 2]; qansel_cnot(&stateVector, qubitCount, a0, a1); - break; + break; case QANSEL_INSTRUCTION_SWAP: a0 = program[PC + 1]; a1 = program[PC + 2]; qansel_swap(&stateVector, qubitCount, a0, a1); - break; + break; case QANSEL_INSTRUCTION_CCX: a0 = program[PC + 1]; a1 = program[PC + 2]; a2 = program[PC + 3]; qansel_toffoli(&stateVector, qubitCount, a0, a1, a2); - break; + break; case QANSEL_INSTRUCTION_CSWAP: a0 = program[PC + 1]; a1 = program[PC + 2]; a2 = program[PC + 3]; qansel_fredkin(&stateVector, qubitCount, a0, a1, a2); - break; + break; case QANSEL_INSTRUCTION_MEASURE: a0 = program[PC + 1]; - a1 = program[PC + 2] - 0x10; + a1 = program[PC + 2] - QANSEL_CBOUND_LOWER; bitVector[a1] = qansel_measure(&stateVector, qubitCount, a0); - break; + break; case QANSEL_INSTRUCTION_BORN: a0 = program[PC + 1]; qansel_born(&stateVector, PC, qubitCount, a0); - break; + break; case QANSEL_INSTRUCTION_DENSITY: a0 = program[PC + 1]; qansel_density_or_print(&stateVector, bitVector, 1, bitCount, qubitCount, a0); - break; + break; case QANSEL_INSTRUCTION_PRINT: a0 = program[PC + 1]; qansel_density_or_print(&stateVector, bitVector, 0, bitCount, qubitCount, a0); - break; + break; case QANSEL_INSTRUCTION_BARRIER: a0 = program[PC + 1]; - break; + break; case QANSEL_INSTRUCTION_RESET: a0 = program[PC + 1]; qansel_reset(&stateVector, bitVector, qubitCount, bitCount, a0); - break; + break; case QANSEL_INSTRUCTION_HVAR: QANSEL_HIDDEN_VARIABLE = 1; float tmp1 = qansel_get_float(program, PC + 1); unsigned int tmp2; memcpy(&tmp2, &tmp1, sizeof(unsigned int)); srand(tmp2); - break; + break; case QANSEL_INSTRUCTION_RAND: QANSEL_HIDDEN_VARIABLE = 0; - break; + break; case QANSEL_INSTRUCTION_IF_E: a0 = program[PC + 1]; tmp = qansel_get_short(program, PC + 2); flags = qansel_compare(bitVector, bitCount, PC, a0, tmp); skip = 1; if (flags & QANSEL_FLAGS_EQUAL) skip = 0; - break; + break; case QANSEL_INSTRUCTION_IF_NE: a0 = program[PC + 1]; tmp = qansel_get_short(program, PC + 2); flags = qansel_compare(bitVector, bitCount, PC, a0, tmp); skip = 1; if (!(flags & QANSEL_FLAGS_EQUAL)) skip = 0; - break; + break; case QANSEL_INSTRUCTION_IF_G: a0 = program[PC + 1]; tmp = qansel_get_short(program, PC + 2); flags = qansel_compare(bitVector, bitCount, PC, a0, tmp); skip = 1; if (flags & QANSEL_FLAGS_GREATER) skip = 0; - break; + break; case QANSEL_INSTRUCTION_IF_GE: a0 = program[PC + 1]; tmp = qansel_get_short(program, PC + 2); flags = qansel_compare(bitVector, bitCount, PC, a0, tmp); skip = 1; if ((flags & QANSEL_FLAGS_GREATER) && (flags & QANSEL_FLAGS_EQUAL)) skip = 0; - break; + break; case QANSEL_INSTRUCTION_IF_L: a0 = program[PC + 1]; tmp = qansel_get_short(program, PC + 2); flags = qansel_compare(bitVector, bitCount, PC, a0, tmp); skip = 1; if (flags & QANSEL_FLAGS_LESSER) skip = 0; - break; + break; case QANSEL_INSTRUCTION_IF_LE: a0 = program[PC + 1]; tmp = qansel_get_short(program, PC + 2); flags = qansel_compare(bitVector, bitCount, PC, a0, tmp); skip = 1; if ((flags & QANSEL_FLAGS_LESSER) && (flags & QANSEL_FLAGS_EQUAL)) skip = 0; - break; + break; case QANSEL_INSTRUCTION_SAMPLE: break; } @@ -1084,7 +1112,7 @@ void qansel_run(unsigned char* program, int programSize, int qubitCount, int bit { for (int i = 0; i < qubitCount; i++) { - cpx_mtx_free(&(queueVector[i])); + free(queueVector[i].ptr); } free(queueVector); } @@ -1109,7 +1137,7 @@ int qanselExecuteBytecode(unsigned char* buff, int sizeofbuff) return 0; } - if (sample != 0xFF) + if (sample != QANSEL_ALL) { unsigned short stats[65536]; for (unsigned int i = 0; i < (1 << bitCount); i++) @@ -1141,18 +1169,18 @@ int qanselExecuteBytecode(unsigned char* buff, int sizeofbuff) { count += stats[i]; } - if (sample == 0x0F) + if (sample == QANSEL_ALL_QUANTUM) { putchar('0' + bit); } tmp <<= 1; } - if (sample == 0x0F) + if (sample == QANSEL_ALL_QUANTUM) { printf(": %.1f%%\n", ((float)stats[i] / (float)shots) * (float)100); } } - if (sample != 0x0F) + if (sample != QANSEL_ALL_QUANTUM) { float prob = ((float)count / (float)shots) * (float)100; printf("0: %.1f%%\n", ((float)100)-prob); diff --git a/src/kernel.cl b/src/kernel.cl index d11e287..dcfdf50 100644 --- a/src/kernel.cl +++ b/src/kernel.cl @@ -20,17 +20,13 @@ __kernel void kernel_dot float rR = 0; float iR = 0; - const int posR = rowR * (colsR * 2) + (colR * 2); for (int i = 0; i < colsA; i++) { - const int posA = rowR * (colsA * 2) + (i * 2); - const int posB = i * (colsB * 2) + (colR * 2); - - const float rA = ptrA[posA]; - const float iA = ptrA[posA + 1]; - const float rB = ptrB[posB]; - const float iB = ptrB[posB + 1]; + 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; @@ -41,8 +37,8 @@ __kernel void kernel_dot rR += first + lasts; iR += outer + inner; } - ptrR[rowR * (colsR * 2) + (colR * 2)] = rR; - ptrR[rowR * (colsR * 2) + (colR * 2) + 1] = iR; + 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 @@ -112,9 +108,8 @@ __kernel void kernel_knk_2x2 { const int rowA = rowR / 2; const int colA = colR / 2; - const int posA = rowA * (colsA * 2) + (colA * 2); - const float rA = ptrA[posA]; - const float iA = ptrA[posA + 1]; + 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++) { @@ -135,20 +130,20 @@ __kernel void kernel_knk_2x2 switch (i) { case 0: - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts; - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner; + 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[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts; - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner; + 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[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts; - ptrR[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner; + 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[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts; - ptrR[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner; + 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 index f8a12f2..b0087ae 100644 --- a/src/kernel_cpu.cl +++ b/src/kernel_cpu.cl @@ -18,17 +18,13 @@ void kernel_dot float rR = 0; float iR = 0; - const int posR = rowR * (colsR * 2) + (colR * 2); for (int i = 0; i < colsA; i++) { - const int posA = rowR * (colsA * 2) + (i * 2); - const int posB = i * (colsB * 2) + (colR * 2); - - const float rA = ptrA[posA]; - const float iA = ptrA[posA + 1]; - const float rB = ptrB[posB]; - const float iB = ptrB[posB + 1]; + 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; @@ -39,8 +35,8 @@ void kernel_dot rR += first + lasts; iR += outer + inner; } - ptrR[rowR * (colsR * 2) + (colR * 2)] = rR; - ptrR[rowR * (colsR * 2) + (colR * 2) + 1] = iR; + 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 @@ -108,9 +104,8 @@ void kernel_knk_2x2 { const int rowA = rowR / 2; const int colA = colR / 2; - const int posA = rowA * (colsA * 2) + (colA * 2); - const float rA = ptrA[posA]; - const float iA = ptrA[posA + 1]; + 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++) { @@ -131,20 +126,20 @@ void kernel_knk_2x2 switch (i) { case 0: - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts; - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner; + 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[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts; - ptrR[(rowR + 0) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner; + 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[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2)] = first + lasts; - ptrR[(rowR + 1) * (colsR * 2) + ((colR + 0) * 2) + 1] = outer + inner; + 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[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2)] = first + lasts; - ptrR[(rowR + 1) * (colsR * 2) + ((colR + 1) * 2) + 1] = outer + inner; + 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_gpu.cl b/src/kernel_gpu.cl index e5e8787..47bf476 100644 --- a/src/kernel_gpu.cl +++ b/src/kernel_gpu.cl @@ -31,342 +31,406 @@ unsigned char kernel_gpu[] = { 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, 0x20, - 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x69, 0x6e, 0x74, - 0x20, 0x70, 0x6f, 0x73, 0x52, 0x20, 0x3d, 0x20, 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, 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, 0x69, 0x6e, 0x74, 0x20, 0x70, 0x6f, 0x73, 0x41, 0x20, 0x3d, - 0x20, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, - 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x69, - 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, 0x69, 0x20, 0x2a, 0x20, - 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, - 0x2b, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x52, 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, + 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, - 0x69, 0x41, 0x20, 0x3d, 0x20, 0x70, 0x74, 0x72, 0x41, 0x5b, 0x70, 0x6f, - 0x73, 0x41, 0x20, 0x2b, 0x20, 0x31, 0x5d, 0x3b, 0x0a, 0x20, 0x20, 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, 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, + 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, 0x66, 0x69, 0x72, 0x73, 0x74, 0x20, 0x3d, - 0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x72, 0x42, 0x3b, 0x0a, 0x20, 0x20, + 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, 0x6f, 0x75, 0x74, 0x65, 0x72, 0x20, - 0x3d, 0x20, 0x72, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, 0x3b, 0x0a, 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, 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, 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, 0x72, 0x52, 0x3b, 0x0a, 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, 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, + 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, 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, + 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, 0x63, 0x6f, 0x6c, 0x41, 0x20, - 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2f, 0x20, 0x63, 0x6f, 0x6c, + 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, 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, + 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, 0x63, 0x6f, 0x6c, 0x42, 0x20, 0x3d, 0x20, 0x63, 0x6f, 0x6c, 0x52, - 0x20, 0x25, 0x20, 0x63, 0x6f, 0x6c, 0x73, 0x42, 0x3b, 0x0a, 0x0a, 0x20, + 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, 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, 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, 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, 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, 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, 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, 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, + 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, 0x30, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, + 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, 0x31, 0x2c, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x63, 0x6f, 0x6e, 0x73, + 0x65, 0x34, 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, + 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, 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, + 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, 0x72, 0x6f, 0x77, 0x73, 0x52, 0x20, 0x3d, 0x20, 0x72, - 0x6f, 0x77, 0x73, 0x41, 0x20, 0x2a, 0x20, 0x32, 0x3b, 0x0a, 0x20, 0x20, + 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, - 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, 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, 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, + 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, 0x70, 0x6f, 0x73, 0x41, 0x20, 0x2b, 0x20, 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, + 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, 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, 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, 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, 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, - 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, + 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, 0x6c, 0x61, 0x73, - 0x74, 0x73, 0x20, 0x3d, 0x20, 0x69, 0x41, 0x20, 0x2a, 0x20, 0x69, 0x42, + 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, 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, 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, - 0x30, 0x3a, 0x0a, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 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, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, - 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, - 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, - 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29, 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, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, - 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x72, 0x6f, 0x77, 0x52, - 0x20, 0x2b, 0x20, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, - 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, - 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29, 0x20, 0x2a, 0x20, - 0x32, 0x29, 0x20, 0x2b, 0x20, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, + 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, 0x31, 0x3a, 0x0a, + 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, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29, 0x20, - 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, - 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, - 0x20, 0x31, 0x29, 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, 0x20, - 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, - 0x74, 0x72, 0x52, 0x5b, 0x28, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, - 0x30, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, - 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, - 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, - 0x2b, 0x20, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x6f, 0x75, 0x74, 0x65, 0x72, + 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, 0x32, 0x3a, 0x0a, 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, 0x72, - 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, - 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, - 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x30, 0x29, - 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, 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, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, - 0x5b, 0x28, 0x72, 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, - 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, - 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, - 0x20, 0x30, 0x29, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x31, + 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, 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, 0x72, 0x6f, 0x77, 0x52, - 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, 0x63, 0x6f, 0x6c, - 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, 0x20, 0x28, 0x28, - 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 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, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, 0x20, - 0x20, 0x20, 0x20, 0x20, 0x20, 0x70, 0x74, 0x72, 0x52, 0x5b, 0x28, 0x72, - 0x6f, 0x77, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2a, 0x20, 0x28, - 0x63, 0x6f, 0x6c, 0x73, 0x52, 0x20, 0x2a, 0x20, 0x32, 0x29, 0x20, 0x2b, - 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6c, 0x52, 0x20, 0x2b, 0x20, 0x31, 0x29, - 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, 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 + 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 }; -unsigned int kernel_gpu_len = 4428; +unsigned int kernel_gpu_len = 5187; diff --git a/src/oldex.mm b/src/oldex.mm deleted file mode 100644 index bbb35ee..0000000 --- a/src/oldex.mm +++ /dev/null @@ -1,127 +0,0 @@ - - //0+0 - QANSEL_INSTRUCTION_CCX, 0x00, 0x01, 0x02, - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x10, - QANSEL_INSTRUCTION_MEASURE, 0x02, 0x11, - QANSEL_INSTRUCTION_PRINT, 0x1F, - QANSEL_INSTRUCTION_RESET, 0xFF, - - //0+1 - QANSEL_INSTRUCTION_X, 0x00, - QANSEL_INSTRUCTION_CCX, 0x00, 0x01, 0x02, - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x10, - QANSEL_INSTRUCTION_MEASURE, 0x02, 0x11, - QANSEL_INSTRUCTION_PRINT, 0x1F, - QANSEL_INSTRUCTION_RESET, 0xFF, - - //1+0 - QANSEL_INSTRUCTION_X, 0x01, - QANSEL_INSTRUCTION_CCX, 0x00, 0x01, 0x02, - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x10, - QANSEL_INSTRUCTION_MEASURE, 0x02, 0x11, - QANSEL_INSTRUCTION_PRINT, 0x1F, - QANSEL_INSTRUCTION_RESET, 0xFF, - - //1+1 - QANSEL_INSTRUCTION_X, 0x00, - QANSEL_INSTRUCTION_X, 0x01, - QANSEL_INSTRUCTION_CCX, 0x00, 0x01, 0x02, - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x10, - QANSEL_INSTRUCTION_MEASURE, 0x02, 0x11, - QANSEL_INSTRUCTION_PRINT, 0x1F, - QANSEL_INSTRUCTION_RESET, 0xFF, - - - - - QANSEL_INSTRUCTION_U3, 0x00, 0xED, 00, 00, 00, 0xED, 00, 00, 00, 0xED, 00, 00, 00, - QANSEL_INSTRUCTION_PRINT, 0x00, - - QANSEL_INSTRUCTION_H, 0x01, - QANSEL_INSTRUCTION_CX, 0x01, 0x02, - - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - QANSEL_INSTRUCTION_H, 0x00, - - QANSEL_INSTRUCTION_MEASURE, 0x00, 0x10, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x11, - - QANSEL_INSTRUCTION_IF_E, 0x1F, 0xED, 0x00, - QANSEL_INSTRUCTION_Z, 0x02, - QANSEL_INSTRUCTION_IF_E, 0x1F, 0xED, 0x00, - QANSEL_INSTRUCTION_X, 0x02, - QANSEL_INSTRUCTION_IF_E, 0x1F, 0xED, 0x00, - QANSEL_INSTRUCTION_X, 0x02, - QANSEL_INSTRUCTION_IF_E, 0x1F, 0xED, 0x00, - QANSEL_INSTRUCTION_Z, 0x02, - - QANSEL_INSTRUCTION_RESET, 0x00, - QANSEL_INSTRUCTION_RESET, 0x01, - QANSEL_INSTRUCTION_PRINT, 0x02, - - - - //Host generates random bits - // for the two players. - QANSEL_INSTRUCTION_H, 0x00, - QANSEL_INSTRUCTION_H, 0x01, - QANSEL_INSTRUCTION_MEASURE, 0x00, 0x10, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x11, - QANSEL_INSTRUCTION_RESET, 0x00, - QANSEL_INSTRUCTION_RESET, 0x01, - - //Two players are also provided - // an entangled qubit. - QANSEL_INSTRUCTION_H, 0x00, - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - - //Player X strategy - //if(c[0]==0) do nothing - QANSEL_INSTRUCTION_IF_E, 0x10, 0xEE, 0x00, - QANSEL_INSTRUCTION_RY, 0x00, 0xEE, 0x00, 0x00, 0x00, - - //Player Y strategy - QANSEL_INSTRUCTION_IF_E, 0x11, 0xEE, 0x00, - QANSEL_INSTRUCTION_RY, 0x01, 0xEE, 0x00, 0x00, 0x00, - QANSEL_INSTRUCTION_IF_E, 0x11, 0xEE, 0x00, - QANSEL_INSTRUCTION_RY, 0x01, 0xEE, 0x00, 0x00, 0x00, - - //Transfer to host - QANSEL_INSTRUCTION_MEASURE, 0x00, 0x12, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x13, - - //a xor b - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x13, - - //Load x and y - QANSEL_INSTRUCTION_RESET, 0x00, - QANSEL_INSTRUCTION_IF_E, 0x10, 0xEE, 0x00, - QANSEL_INSTRUCTION_X, 0x00, - QANSEL_INSTRUCTION_RESET, 0x01, - QANSEL_INSTRUCTION_IF_E, 0x11, 0xEE, 0x00, - QANSEL_INSTRUCTION_X, 0x01, - - //x and y - QANSEL_INSTRUCTION_CCX, 0x00, 0x01, 0x02, - QANSEL_INSTRUCTION_MEASURE, 0x02, 0x12, - - //Load (a xor b) and (x and y) - QANSEL_INSTRUCTION_RESET, 0x00, - QANSEL_INSTRUCTION_IF_E, 0x12, 0xEE, 0x00, - QANSEL_INSTRUCTION_X, 0x00, - QANSEL_INSTRUCTION_RESET, 0x01, - QANSEL_INSTRUCTION_IF_E, 0x13, 0xEE, 0x00, - QANSEL_INSTRUCTION_X, 0x01, - - //(a xor b) = (x and y) - QANSEL_INSTRUCTION_CX, 0x00, 0x01, - QANSEL_INSTRUCTION_X, 0x01, - - //Store final results - QANSEL_INSTRUCTION_MEASURE, 0x01, 0x10, - QANSEL_INSTRUCTION_SAMPLE, 0x10 \ No newline at end of file diff --git a/src/openqasm.c b/src/openqasm.c index f451fa5..2a3fbaa 100644 --- a/src/openqasm.c +++ b/src/openqasm.c @@ -320,7 +320,7 @@ int qansel_process_chunk(int index, char* chunk, int line, regmatch_t* regmatche } else { - a0 = 0x0F; + a0 = QANSEL_ALL_QUANTUM; if (instr == QANSEL_INSTRUCTION_DENSITY) { fprintf(stderr, "QAnsel on line %i: Density matrices can only be produced for individual qubits.\n", line, a0); @@ -396,7 +396,7 @@ int qansel_process_chunk(int index, char* chunk, int line, regmatch_t* regmatche } else { - a0 = 0x0F; + a0 = QANSEL_ALL_QUANTUM; } } { @@ -417,11 +417,11 @@ int qansel_process_chunk(int index, char* chunk, int line, regmatch_t* regmatche fprintf(stderr, "QAnsel on line %i: Index `%i` exceeds allocated classical bits.\n", line, a1); return 0; } - a1 += 0x10; + a1 += QANSEL_CBOUND_LOWER; } else { - a1 = 0x1F; + a1 = QANSEL_ALL_CLASSIC; } } *binarySize += 3; @@ -467,11 +467,11 @@ int qansel_process_chunk(int index, char* chunk, int line, regmatch_t* regmatche fprintf(stderr, "QAnsel on line %i: Index `%i` exceeds allocated classical bits.\n", line, a0); return 0; } - a0 += 0x10; + a0 += QANSEL_CBOUND_LOWER; } else { - a0 = 0x1F; + a0 = QANSEL_ALL_CLASSIC; } } *binarySize += 2; @@ -567,11 +567,11 @@ int qansel_process_chunk(int index, char* chunk, int line, regmatch_t* regmatche fprintf(stderr, "QAnsel on line %i: Index `%i` exceeds allocated classical bits.\n", line, a0); return 0; } - a0 += 0x10; + a0 += QANSEL_CBOUND_LOWER; } else { - a0 = 0x1F; + a0 = QANSEL_ALL_CLASSIC; } } { diff --git a/src/qansel.h b/src/qansel.h index ef55269..37cb7a6 100644 --- a/src/qansel.h +++ b/src/qansel.h @@ -11,16 +11,17 @@ #include "gates.c" #include "display.c" #include "chacha20.c" -#define QANSEL_QUBITS_MAX 14 +#define QANSEL_QUBITS_MAX 16 unsigned char QANSEL_QUBIT_LIMIT = QANSEL_QUBITS_MAX; unsigned char QANSEL_HIDDEN_VARIABLE = 0; unsigned char QANSEL_USE_DISPLAY = 0; unsigned char QANSEL_VERBOSE = 0; +unsigned char QANSEL_USE_QUEUE = 0; FILE* QANSEL_RANDOM_FILE = NULL; #define QANSEL_MODE_BARE 0 #define QANSEL_MODE_THREADED 1 #define QANSEL_MODE_METAL 2 -#define QANSEL_MODE_METAL_THREADED 3 +#define QANSEL_MODE_QUEUE 4 unsigned char QANSEL_MODE = QANSEL_MODE_BARE; //#define SPEED_TEST @@ -61,4 +62,14 @@ unsigned char QANSEL_MODE = QANSEL_MODE_BARE; #define QANSEL_FLAGS_GREATER 0b00010 #define QANSEL_FLAGS_LESSER 0b00100 +#define QANSEL_QBOUND_LOWER 0x00 +#define QANSEL_QBOUND_UPPER 0x0F +#define QANSEL_CBOUND_LOWER 0x10 +#define QANSEL_CBOUND_UPPER 0x1F + +#define QANSEL_ALL_QUANTUM 0x2F +#define QANSEL_ALL_CLASSIC 0x3F +#define QANSEL_ALL 0xFF + + #endif \ No newline at end of file -- 2.39.5