From 65f8f7c5280898a4663ed0c37a9613aa634fa72a Mon Sep 17 00:00:00 2001 From: "Dr.Elias" Date: Thu, 30 Dec 2021 00:09:00 +0300 Subject: [PATCH] opencl optimisations --- crypto/util/cruncher.h | 342 ++++++++++++++++++++++++++++++++++ crypto/util/cruncher_h.h | 2 + crypto/util/opencl/opencl.cpp | 135 +++++--------- crypto/util/opencl/opencl.h | 21 ++- crypto/util/opencl/sha256.cpp | 17 +- 5 files changed, 410 insertions(+), 107 deletions(-) create mode 100644 crypto/util/cruncher.h create mode 100644 crypto/util/cruncher_h.h diff --git a/crypto/util/cruncher.h b/crypto/util/cruncher.h new file mode 100644 index 00000000..c314bb1d --- /dev/null +++ b/crypto/util/cruncher.h @@ -0,0 +1,342 @@ +#ifndef __CRUNCHER_H__ +#define __CRUNCHER_H__ + + +#define MAX_GPU_THREADS 16 + + +#if defined(__OPENCL_VERSION__) + typedef uint uint32_t; + typedef ulong uint64_t; +#else + #if defined(__NVCC__) + #include + #include + #endif + + #include +#endif + + +#define ROTRIGHT(a, b) (((a) >> (b)) | ((a) << (32 - (b)))) + +#define CH(x, y, z) (((x) & (y)) ^ (~(x) & (z))) +#define MAJ(x, y, z) (((x) & (y)) ^ ((x) & (z)) ^ ((y) & (z))) +#define EP0(x) (ROTRIGHT(x, 2) ^ ROTRIGHT(x, 13) ^ ROTRIGHT(x, 22)) +#define EP1(x) (ROTRIGHT(x, 6) ^ ROTRIGHT(x, 11) ^ ROTRIGHT(x, 25)) +#define SIG0(x) (ROTRIGHT(x, 7) ^ ROTRIGHT(x, 18) ^ ((x) >> 3)) +#define SIG1(x) (ROTRIGHT(x, 17) ^ ROTRIGHT(x, 19) ^ ((x) >> 10)) + + +typedef struct vec8u +{ + uint32_t v[8]; +} +vec8u; + + +typedef struct vec16u +{ + uint32_t v[16]; +} +vec16u; + + +typedef struct ThreadData +{ + vec8u state; + uint32_t rdata[9]; +} +ThreadData; + + +typedef struct MsgData +{ + uint32_t pseed[3]; + ThreadData thrdata[MAX_GPU_THREADS]; + vec8u target; +} +MsgData; + + +#if defined(__NVCC__) +__forceinline__ __device__ __host__ +#else +inline +#endif +static vec8u sha256_transform(vec16u data, vec8u state) +{ + const uint32_t ksha[] = { + 0x428A2F98, 0x71374491, 0xB5C0FBCF, 0xE9B5DBA5, 0x3956C25B, 0x59F111F1, 0x923F82A4, 0xAB1C5ED5, + 0xD807AA98, 0x12835B01, 0x243185BE, 0x550C7DC3, 0x72BE5D74, 0x80DEB1FE, 0x9BDC06A7, 0xC19BF174, + 0xE49B69C1, 0xEFBE4786, 0x0FC19DC6, 0x240CA1CC, 0x2DE92C6F, 0x4A7484AA, 0x5CB0A9DC, 0x76F988DA, + 0x983E5152, 0xA831C66D, 0xB00327C8, 0xBF597FC7, 0xC6E00BF3, 0xD5A79147, 0x06CA6351, 0x14292967, + 0x27B70A85, 0x2E1B2138, 0x4D2C6DFC, 0x53380D13, 0x650A7354, 0x766A0ABB, 0x81C2C92E, 0x92722C85, + 0xA2BFE8A1, 0xA81A664B, 0xC24B8B70, 0xC76C51A3, 0xD192E819, 0xD6990624, 0xF40E3585, 0x106AA070, + 0x19A4C116, 0x1E376C08, 0x2748774C, 0x34B0BCB5, 0x391C0CB3, 0x4ED8AA4A, 0x5B9CCA4F, 0x682E6FF3, + 0x748F82EE, 0x78A5636F, 0x84C87814, 0x8CC70208, 0x90BEFFFA, 0xA4506CEB, 0xBEF9A3F7, 0xC67178F2 + }; + + uint32_t m[64]; + + for (int i = 0; i < 16; i += 1) + m[i] = data.v[i]; + + for (int i = 16; i < 64; i += 1) + m[i] = SIG1(m[i - 2]) + m[i - 7] + SIG0(m[i - 15]) + m[i - 16]; + + uint32_t a = state.v[0], b = state.v[1], c = state.v[2], d = state.v[3]; + uint32_t e = state.v[4], f = state.v[5], g = state.v[6], h = state.v[7]; + +#if defined(__CUDA_ARCH__) || defined(__OPENCL_VERSION__) + #pragma unroll +#endif + for (int i = 0; i < 64; i += 1) + { + uint32_t t1 = h + EP1(e) + CH(e, f, g) + ksha[i] + m[i]; + uint32_t t2 = EP0(a) + MAJ(a, b, c); + h = g; + g = f; + f = e; + e = d + t1; + d = c; + c = b; + b = a; + a = t1 + t2; + } + + state.v[0] += a, state.v[1] += b, state.v[2] += c, state.v[3] += d; + state.v[4] += e, state.v[5] += f, state.v[6] += g, state.v[7] += h; + + return state; +} + + +#if !defined(__OPENCL_VERSION__) +static const vec8u h256 = { + 0x6A09E667, 0xBB67AE85, 0x3C6EF372, 0xA54FF53A, 0x510E527F, 0x9B05688C, 0x1F83D9AB, 0x5BE0CD19 +}; + + +inline static void bitcredit_be32enc(void *pp, uint32_t x) +{ + uint8_t *p = (uint8_t *)pp; + + p[3] = x & 0xff; + p[2] = (x >> 8) & 0xff; + p[1] = (x >> 16) & 0xff; + p[0] = (x >> 24) & 0xff; +} + + +inline static MsgData bitcredit_prepare_msg(uint32_t gpu_threads, uint32_t expired, const unsigned char *data, + const void *target, const unsigned char *rdata) +{ + /*printf("gpu_threads: %u, expired: %04x\n", (unsigned)gpu_threads, (unsigned)expired); + printf("data:\n"); + + unsigned char padded_data[128]; + memset(padded_data, 0, sizeof(padded_data)); + memcpy(padded_data, data, 123); + + for (int i = 0; i < 32; i += 1) + { + for (int j = 0; j < 4; j += 1) + printf("%02hhx ", padded_data[i * 4 + j]); + printf("\n"); + } + + printf("target:\n"); + for (int i = 0; i < 32; i += 1) + printf("%02hhx ", ((const unsigned char *)target)[i]); + printf("\n"); + + printf("rdata:\n"); + for (unsigned i = 0; i < gpu_threads; i += 1) + { + for (int j = 0; j < 32; j += 1) + printf("%02hhx ", rdata[i * 32 + j]); + printf("\n"); + }*/ + + MsgData msg; + + for (int i = 0; i < 3; i += 1) + bitcredit_be32enc(&msg.pseed[i], *((const uint32_t *)(data + 76) + i)); + + for (int i = 0; i < 8; i += 1) + bitcredit_be32enc(&msg.target.v[i], *((const uint32_t *)target + i)); + + for (uint32_t gpu_thread = 0; gpu_thread < gpu_threads; gpu_thread += 1) + { + ThreadData *thrdata = &msg.thrdata[gpu_thread]; + const unsigned char *thr_rdata = rdata + gpu_thread * 32; + + vec8u state = h256; + vec16u shadata; + + bitcredit_be32enc(&shadata.v[0x0], *((const uint32_t *)data + 0x0)); + + uint32_t word; + bitcredit_be32enc(&word, *((const uint32_t *)data + 0x1)); + + shadata.v[0x1] = (word & 0xFFFFFF00) | (expired >> 24); + shadata.v[0x2] = (expired << 8) | data[11]; + + for (int i = 0x3; i <= 0x9; i += 1) + bitcredit_be32enc(&shadata.v[i], *((const uint32_t *)data + i)); + + bitcredit_be32enc(&word, *((const uint32_t *)data + 0xA)); + shadata.v[0xA] = (word & 0xFFFFFF00) | thr_rdata[0]; + + for (int i = 0xB, j = 0; i <= 0xF; i += 1, j += 1) + bitcredit_be32enc(&shadata.v[i], *((const uint32_t *)(thr_rdata + 1) + j)); + + thrdata->state = sha256_transform(shadata, state); + + for (int i = 0x0, j = 0; i <= 0x1; i += 1, j += 1) + bitcredit_be32enc(&thrdata->rdata[i], *((const uint32_t *)(thr_rdata + 21) + j)); + + bitcredit_be32enc(&word, *(const uint32_t *)(thr_rdata + 28)); + thrdata->rdata[0x2] = (word << 8) | data[75]; + + bitcredit_be32enc(&word, *((const uint32_t *)data + 0x16)); + thrdata->rdata[0x3] = (word & 0xFFFFFF00) | thr_rdata[0]; + + for (int i = 0x4, j = 0; i <= 0x8; i += 1, j += 1) + bitcredit_be32enc(&thrdata->rdata[i], *((const uint32_t *)(thr_rdata + 1) + j)); + } + + return msg; +} +#endif + + +typedef struct DevHashResult +{ + uint64_t nonce; + uint32_t vcpu, found; +} +DevHashResult; + + +#if defined(__CUDACC__) || defined(__OPENCL_VERSION__) + +#if defined(__OPENCL_VERSION__) + __kernel void bitcredit_gpu_hash( + uint64_t start_nonce, __constant const MsgData *c_msg, __global DevHashResult *result + ) +#else + __constant__ MsgData c_msg; + extern "C" __global__ void bitcredit_gpu_hash(uint64_t start_nonce, DevHashResult *result) +#endif +{ +#if defined(__OPENCL_VERSION__) + uint64_t idx = get_global_id(0); + uint32_t vcpu = get_global_id(1); + + __constant const MsgData *msg = c_msg; + __constant const ThreadData *thrdata = &msg->thrdata[vcpu]; +#else + uint64_t idx = blockIdx.x * blockDim.x + threadIdx.x; + uint32_t vcpu = blockIdx.y; + + const MsgData *msg = &c_msg; + const ThreadData *thrdata = &c_msg.thrdata[vcpu]; +#endif + + vec8u state = thrdata->state; + + vec16u data; + uint64_t nonce = start_nonce + idx; + + uint32_t rdata6 = thrdata->rdata[0], rdata7 = thrdata->rdata[1], rdata8 = thrdata->rdata[2]; + uint64_t rdata = (((uint64_t)rdata6 << 56) | ((uint64_t)rdata7 << 24) | ((uint64_t)rdata8 >> 8)) + nonce; + + rdata6 = (uint32_t)(rdata >> 56) | (rdata6 & 0xFFFFFF00); + rdata7 = (uint32_t)(rdata >> 24); + + uint32_t rdata10 = (uint32_t)(rdata << 8) | 0x80; + rdata8 = (uint32_t)(rdata << 8) | (rdata8 & 0xFF); + + data.v[0x0] = rdata6; + data.v[0x1] = rdata7; + data.v[0x2] = rdata8; + data.v[0x3] = msg->pseed[0]; + data.v[0x4] = msg->pseed[1]; + data.v[0x5] = msg->pseed[2]; + data.v[0x6] = thrdata->rdata[3]; + data.v[0x7] = thrdata->rdata[4]; + data.v[0x8] = thrdata->rdata[5]; + data.v[0x9] = thrdata->rdata[6]; + data.v[0xA] = thrdata->rdata[7]; + data.v[0xB] = thrdata->rdata[8]; + data.v[0xC] = rdata6; + data.v[0xD] = rdata7; + data.v[0xE] = rdata10; + data.v[0xF] = 0x00000000; + + /*if (vcpu == 0 && idx == 0) + printf( + "%04x %04x %04x %04x %04x %04x %04x %04x\n%04x %04x %04x %04x %04x %04x %04x %04x\n\n", + data.v[0], data.v[1], data.v[2], data.v[3], data.v[4], data.v[5], data.v[6], data.v[7], + data.v[8], data.v[9], data.v[10], data.v[11], data.v[12], data.v[13], data.v[14], data.v[15] + );*/ + + state = sha256_transform(data, state); + + data.v[0x0] = 0x00000000; + data.v[0x1] = 0x00000000; + data.v[0x2] = 0x00000000; + data.v[0x3] = 0x00000000; + data.v[0x4] = 0x00000000; + data.v[0x5] = 0x00000000; + data.v[0x6] = 0x00000000; + data.v[0x7] = 0x00000000; + data.v[0x8] = 0x00000000; + data.v[0x9] = 0x00000000; + data.v[0xA] = 0x00000000; + data.v[0xB] = 0x00000000; + data.v[0xC] = 0x00000000; + data.v[0xD] = 0x00000000; + data.v[0xE] = 0x00000000; + data.v[0xF] = 0x000003d8; + + /*if (vcpu == 0 && idx == 0) + printf( + "%04x %04x %04x %04x %04x %04x %04x %04x\n%04x %04x %04x %04x %04x %04x %04x %04x\n\n", + data.v[0], data.v[1], data.v[2], data.v[3], data.v[4], data.v[5], data.v[6], data.v[7], + data.v[8], data.v[9], data.v[10], data.v[11], data.v[12], data.v[13], data.v[14], data.v[15] + );*/ + + state = sha256_transform(data, state); + + /*if (vcpu == 0 && idx == 0) + printf( + "%04x %04x %04x %04x %04x %04x %04x %04x\n\n", + state.v[0], state.v[1], state.v[2], state.v[3], state.v[4], state.v[5], state.v[6], state.v[7] + );*/ + + for (int i = 0; i < 8; i += 1) + { + if (state.v[i] > msg->target.v[i]) + return; + + if (state.v[i] < msg->target.v[i]) + { +#if defined(__OPENCL_VERSION__) + if (atomic_add(&result->found, 1) == 0) +#else + if (atomicAdd(&result->found, 1) == 0) +#endif + result->nonce = nonce, result->vcpu = vcpu; + + return; + } + } +} + +#endif + + +#endif diff --git a/crypto/util/cruncher_h.h b/crypto/util/cruncher_h.h new file mode 100644 index 00000000..82e6bff1 --- /dev/null +++ b/crypto/util/cruncher_h.h @@ -0,0 +1,2 @@ +unsigned char cruncher_h[] = {0x23, 0x69, 0x66, 0x6e, 0x64, 0x65, 0x66, 0x20, 0x5f, 0x5f, 0x43, 0x52, 0x55, 0x4e, 0x43, 0x48, 0x45, 0x52, 0x5f, 0x48, 0x5f, 0x5f, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x5f, 0x5f, 0x43, 0x52, 0x55, 0x4e, 0x43, 0x48, 0x45, 0x52, 0x5f, 0x48, 0x5f, 0x5f, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x4d, 0x41, 0x58, 0x5f, 0x47, 0x50, 0x55, 0x5f, 0x54, 0x48, 0x52, 0x45, 0x41, 0x44, 0x53, 0x20, 0x31, 0x36, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4f, 0x50, 0x45, 0x4e, 0x43, 0x4c, 0x5f, 0x56, 0x45, 0x52, 0x53, 0x49, 0x4f, 0x4e, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x9, 0x74, 0x79, 0x70, 0x65, 0x64, 0x65, 0x66, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x3b, 0xd, 0xa, 0x9, 0x74, 0x79, 0x70, 0x65, 0x64, 0x65, 0x66, 0x20, 0x75, 0x6c, 0x6f, 0x6e, 0x67, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x3b, 0xd, 0xa, 0x23, 0x65, 0x6c, 0x73, 0x65, 0xd, 0xa, 0x9, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4e, 0x56, 0x43, 0x43, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x9, 0x9, 0x23, 0x69, 0x6e, 0x63, 0x6c, 0x75, 0x64, 0x65, 0x20, 0x3c, 0x63, 0x75, 0x64, 0x61, 0x5f, 0x72, 0x75, 0x6e, 0x74, 0x69, 0x6d, 0x65, 0x2e, 0x68, 0x3e, 0xd, 0xa, 0x9, 0x9, 0x23, 0x69, 0x6e, 0x63, 0x6c, 0x75, 0x64, 0x65, 0x20, 0x3c, 0x64, 0x65, 0x76, 0x69, 0x63, 0x65, 0x5f, 0x6c, 0x61, 0x75, 0x6e, 0x63, 0x68, 0x5f, 0x70, 0x61, 0x72, 0x61, 0x6d, 0x65, 0x74, 0x65, 0x72, 0x73, 0x2e, 0x68, 0x3e, 0xd, 0xa, 0x9, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0xd, 0xa, 0x9, 0x23, 0x69, 0x6e, 0x63, 0x6c, 0x75, 0x64, 0x65, 0x20, 0x3c, 0x73, 0x74, 0x64, 0x69, 0x6e, 0x74, 0x2e, 0x68, 0x3e, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x61, 0x2c, 0x20, 0x62, 0x29, 0x20, 0x28, 0x28, 0x28, 0x61, 0x29, 0x20, 0x3e, 0x3e, 0x20, 0x28, 0x62, 0x29, 0x29, 0x20, 0x7c, 0x20, 0x28, 0x28, 0x61, 0x29, 0x20, 0x3c, 0x3c, 0x20, 0x28, 0x33, 0x32, 0x20, 0x2d, 0x20, 0x28, 0x62, 0x29, 0x29, 0x29, 0x29, 0xd, 0xa, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x43, 0x48, 0x28, 0x78, 0x2c, 0x20, 0x79, 0x2c, 0x20, 0x7a, 0x29, 0x20, 0x28, 0x28, 0x28, 0x78, 0x29, 0x20, 0x26, 0x20, 0x28, 0x79, 0x29, 0x29, 0x20, 0x5e, 0x20, 0x28, 0x7e, 0x28, 0x78, 0x29, 0x20, 0x26, 0x20, 0x28, 0x7a, 0x29, 0x29, 0x29, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x4d, 0x41, 0x4a, 0x28, 0x78, 0x2c, 0x20, 0x79, 0x2c, 0x20, 0x7a, 0x29, 0x20, 0x28, 0x28, 0x28, 0x78, 0x29, 0x20, 0x26, 0x20, 0x28, 0x79, 0x29, 0x29, 0x20, 0x5e, 0x20, 0x28, 0x28, 0x78, 0x29, 0x20, 0x26, 0x20, 0x28, 0x7a, 0x29, 0x29, 0x20, 0x5e, 0x20, 0x28, 0x28, 0x79, 0x29, 0x20, 0x26, 0x20, 0x28, 0x7a, 0x29, 0x29, 0x29, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x45, 0x50, 0x30, 0x28, 0x78, 0x29, 0x20, 0x28, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x32, 0x29, 0x20, 0x5e, 0x20, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x31, 0x33, 0x29, 0x20, 0x5e, 0x20, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x32, 0x32, 0x29, 0x29, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x45, 0x50, 0x31, 0x28, 0x78, 0x29, 0x20, 0x28, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x36, 0x29, 0x20, 0x5e, 0x20, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x31, 0x31, 0x29, 0x20, 0x5e, 0x20, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x32, 0x35, 0x29, 0x29, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x53, 0x49, 0x47, 0x30, 0x28, 0x78, 0x29, 0x20, 0x28, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x37, 0x29, 0x20, 0x5e, 0x20, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x31, 0x38, 0x29, 0x20, 0x5e, 0x20, 0x28, 0x28, 0x78, 0x29, 0x20, 0x3e, 0x3e, 0x20, 0x33, 0x29, 0x29, 0xd, 0xa, 0x23, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x20, 0x53, 0x49, 0x47, 0x31, 0x28, 0x78, 0x29, 0x20, 0x28, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x31, 0x37, 0x29, 0x20, 0x5e, 0x20, 0x52, 0x4f, 0x54, 0x52, 0x49, 0x47, 0x48, 0x54, 0x28, 0x78, 0x2c, 0x20, 0x31, 0x39, 0x29, 0x20, 0x5e, 0x20, 0x28, 0x28, 0x78, 0x29, 0x20, 0x3e, 0x3e, 0x20, 0x31, 0x30, 0x29, 0x29, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x74, 0x79, 0x70, 0x65, 0x64, 0x65, 0x66, 0x20, 0x73, 0x74, 0x72, 0x75, 0x63, 0x74, 0x20, 0x76, 0x65, 0x63, 0x38, 0x75, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x76, 0x5b, 0x38, 0x5d, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0x76, 0x65, 0x63, 0x38, 0x75, 0x3b, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x74, 0x79, 0x70, 0x65, 0x64, 0x65, 0x66, 0x20, 0x73, 0x74, 0x72, 0x75, 0x63, 0x74, 0x20, 0x76, 0x65, 0x63, 0x31, 0x36, 0x75, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x76, 0x5b, 0x31, 0x36, 0x5d, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0x76, 0x65, 0x63, 0x31, 0x36, 0x75, 0x3b, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x74, 0x79, 0x70, 0x65, 0x64, 0x65, 0x66, 0x20, 0x73, 0x74, 0x72, 0x75, 0x63, 0x74, 0x20, 0x54, 0x68, 0x72, 0x65, 0x61, 0x64, 0x44, 0x61, 0x74, 0x61, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x76, 0x65, 0x63, 0x38, 0x75, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x3b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x39, 0x5d, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0x54, 0x68, 0x72, 0x65, 0x61, 0x64, 0x44, 0x61, 0x74, 0x61, 0x3b, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x74, 0x79, 0x70, 0x65, 0x64, 0x65, 0x66, 0x20, 0x73, 0x74, 0x72, 0x75, 0x63, 0x74, 0x20, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x70, 0x73, 0x65, 0x65, 0x64, 0x5b, 0x33, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x54, 0x68, 0x72, 0x65, 0x61, 0x64, 0x44, 0x61, 0x74, 0x61, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x4d, 0x41, 0x58, 0x5f, 0x47, 0x50, 0x55, 0x5f, 0x54, 0x48, 0x52, 0x45, 0x41, 0x44, 0x53, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x76, 0x65, 0x63, 0x38, 0x75, 0x20, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0x3b, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4e, 0x56, 0x43, 0x43, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x5f, 0x5f, 0x66, 0x6f, 0x72, 0x63, 0x65, 0x69, 0x6e, 0x6c, 0x69, 0x6e, 0x65, 0x5f, 0x5f, 0x20, 0x5f, 0x5f, 0x64, 0x65, 0x76, 0x69, 0x63, 0x65, 0x5f, 0x5f, 0x20, 0x5f, 0x5f, 0x68, 0x6f, 0x73, 0x74, 0x5f, 0x5f, 0xd, 0xa, 0x23, 0x65, 0x6c, 0x73, 0x65, 0xd, 0xa, 0x69, 0x6e, 0x6c, 0x69, 0x6e, 0x65, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0x73, 0x74, 0x61, 0x74, 0x69, 0x63, 0x20, 0x76, 0x65, 0x63, 0x38, 0x75, 0x20, 0x73, 0x68, 0x61, 0x32, 0x35, 0x36, 0x5f, 0x74, 0x72, 0x61, 0x6e, 0x73, 0x66, 0x6f, 0x72, 0x6d, 0x28, 0x76, 0x65, 0x63, 0x31, 0x36, 0x75, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2c, 0x20, 0x76, 0x65, 0x63, 0x38, 0x75, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x29, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x6b, 0x73, 0x68, 0x61, 0x5b, 0x5d, 0x20, 0x3d, 0x20, 0x7b, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x34, 0x32, 0x38, 0x41, 0x32, 0x46, 0x39, 0x38, 0x2c, 0x20, 0x30, 0x78, 0x37, 0x31, 0x33, 0x37, 0x34, 0x34, 0x39, 0x31, 0x2c, 0x20, 0x30, 0x78, 0x42, 0x35, 0x43, 0x30, 0x46, 0x42, 0x43, 0x46, 0x2c, 0x20, 0x30, 0x78, 0x45, 0x39, 0x42, 0x35, 0x44, 0x42, 0x41, 0x35, 0x2c, 0x20, 0x30, 0x78, 0x33, 0x39, 0x35, 0x36, 0x43, 0x32, 0x35, 0x42, 0x2c, 0x20, 0x30, 0x78, 0x35, 0x39, 0x46, 0x31, 0x31, 0x31, 0x46, 0x31, 0x2c, 0x20, 0x30, 0x78, 0x39, 0x32, 0x33, 0x46, 0x38, 0x32, 0x41, 0x34, 0x2c, 0x20, 0x30, 0x78, 0x41, 0x42, 0x31, 0x43, 0x35, 0x45, 0x44, 0x35, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x44, 0x38, 0x30, 0x37, 0x41, 0x41, 0x39, 0x38, 0x2c, 0x20, 0x30, 0x78, 0x31, 0x32, 0x38, 0x33, 0x35, 0x42, 0x30, 0x31, 0x2c, 0x20, 0x30, 0x78, 0x32, 0x34, 0x33, 0x31, 0x38, 0x35, 0x42, 0x45, 0x2c, 0x20, 0x30, 0x78, 0x35, 0x35, 0x30, 0x43, 0x37, 0x44, 0x43, 0x33, 0x2c, 0x20, 0x30, 0x78, 0x37, 0x32, 0x42, 0x45, 0x35, 0x44, 0x37, 0x34, 0x2c, 0x20, 0x30, 0x78, 0x38, 0x30, 0x44, 0x45, 0x42, 0x31, 0x46, 0x45, 0x2c, 0x20, 0x30, 0x78, 0x39, 0x42, 0x44, 0x43, 0x30, 0x36, 0x41, 0x37, 0x2c, 0x20, 0x30, 0x78, 0x43, 0x31, 0x39, 0x42, 0x46, 0x31, 0x37, 0x34, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x45, 0x34, 0x39, 0x42, 0x36, 0x39, 0x43, 0x31, 0x2c, 0x20, 0x30, 0x78, 0x45, 0x46, 0x42, 0x45, 0x34, 0x37, 0x38, 0x36, 0x2c, 0x20, 0x30, 0x78, 0x30, 0x46, 0x43, 0x31, 0x39, 0x44, 0x43, 0x36, 0x2c, 0x20, 0x30, 0x78, 0x32, 0x34, 0x30, 0x43, 0x41, 0x31, 0x43, 0x43, 0x2c, 0x20, 0x30, 0x78, 0x32, 0x44, 0x45, 0x39, 0x32, 0x43, 0x36, 0x46, 0x2c, 0x20, 0x30, 0x78, 0x34, 0x41, 0x37, 0x34, 0x38, 0x34, 0x41, 0x41, 0x2c, 0x20, 0x30, 0x78, 0x35, 0x43, 0x42, 0x30, 0x41, 0x39, 0x44, 0x43, 0x2c, 0x20, 0x30, 0x78, 0x37, 0x36, 0x46, 0x39, 0x38, 0x38, 0x44, 0x41, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x39, 0x38, 0x33, 0x45, 0x35, 0x31, 0x35, 0x32, 0x2c, 0x20, 0x30, 0x78, 0x41, 0x38, 0x33, 0x31, 0x43, 0x36, 0x36, 0x44, 0x2c, 0x20, 0x30, 0x78, 0x42, 0x30, 0x30, 0x33, 0x32, 0x37, 0x43, 0x38, 0x2c, 0x20, 0x30, 0x78, 0x42, 0x46, 0x35, 0x39, 0x37, 0x46, 0x43, 0x37, 0x2c, 0x20, 0x30, 0x78, 0x43, 0x36, 0x45, 0x30, 0x30, 0x42, 0x46, 0x33, 0x2c, 0x20, 0x30, 0x78, 0x44, 0x35, 0x41, 0x37, 0x39, 0x31, 0x34, 0x37, 0x2c, 0x20, 0x30, 0x78, 0x30, 0x36, 0x43, 0x41, 0x36, 0x33, 0x35, 0x31, 0x2c, 0x20, 0x30, 0x78, 0x31, 0x34, 0x32, 0x39, 0x32, 0x39, 0x36, 0x37, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x32, 0x37, 0x42, 0x37, 0x30, 0x41, 0x38, 0x35, 0x2c, 0x20, 0x30, 0x78, 0x32, 0x45, 0x31, 0x42, 0x32, 0x31, 0x33, 0x38, 0x2c, 0x20, 0x30, 0x78, 0x34, 0x44, 0x32, 0x43, 0x36, 0x44, 0x46, 0x43, 0x2c, 0x20, 0x30, 0x78, 0x35, 0x33, 0x33, 0x38, 0x30, 0x44, 0x31, 0x33, 0x2c, 0x20, 0x30, 0x78, 0x36, 0x35, 0x30, 0x41, 0x37, 0x33, 0x35, 0x34, 0x2c, 0x20, 0x30, 0x78, 0x37, 0x36, 0x36, 0x41, 0x30, 0x41, 0x42, 0x42, 0x2c, 0x20, 0x30, 0x78, 0x38, 0x31, 0x43, 0x32, 0x43, 0x39, 0x32, 0x45, 0x2c, 0x20, 0x30, 0x78, 0x39, 0x32, 0x37, 0x32, 0x32, 0x43, 0x38, 0x35, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x41, 0x32, 0x42, 0x46, 0x45, 0x38, 0x41, 0x31, 0x2c, 0x20, 0x30, 0x78, 0x41, 0x38, 0x31, 0x41, 0x36, 0x36, 0x34, 0x42, 0x2c, 0x20, 0x30, 0x78, 0x43, 0x32, 0x34, 0x42, 0x38, 0x42, 0x37, 0x30, 0x2c, 0x20, 0x30, 0x78, 0x43, 0x37, 0x36, 0x43, 0x35, 0x31, 0x41, 0x33, 0x2c, 0x20, 0x30, 0x78, 0x44, 0x31, 0x39, 0x32, 0x45, 0x38, 0x31, 0x39, 0x2c, 0x20, 0x30, 0x78, 0x44, 0x36, 0x39, 0x39, 0x30, 0x36, 0x32, 0x34, 0x2c, 0x20, 0x30, 0x78, 0x46, 0x34, 0x30, 0x45, 0x33, 0x35, 0x38, 0x35, 0x2c, 0x20, 0x30, 0x78, 0x31, 0x30, 0x36, 0x41, 0x41, 0x30, 0x37, 0x30, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x31, 0x39, 0x41, 0x34, 0x43, 0x31, 0x31, 0x36, 0x2c, 0x20, 0x30, 0x78, 0x31, 0x45, 0x33, 0x37, 0x36, 0x43, 0x30, 0x38, 0x2c, 0x20, 0x30, 0x78, 0x32, 0x37, 0x34, 0x38, 0x37, 0x37, 0x34, 0x43, 0x2c, 0x20, 0x30, 0x78, 0x33, 0x34, 0x42, 0x30, 0x42, 0x43, 0x42, 0x35, 0x2c, 0x20, 0x30, 0x78, 0x33, 0x39, 0x31, 0x43, 0x30, 0x43, 0x42, 0x33, 0x2c, 0x20, 0x30, 0x78, 0x34, 0x45, 0x44, 0x38, 0x41, 0x41, 0x34, 0x41, 0x2c, 0x20, 0x30, 0x78, 0x35, 0x42, 0x39, 0x43, 0x43, 0x41, 0x34, 0x46, 0x2c, 0x20, 0x30, 0x78, 0x36, 0x38, 0x32, 0x45, 0x36, 0x46, 0x46, 0x33, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x30, 0x78, 0x37, 0x34, 0x38, 0x46, 0x38, 0x32, 0x45, 0x45, 0x2c, 0x20, 0x30, 0x78, 0x37, 0x38, 0x41, 0x35, 0x36, 0x33, 0x36, 0x46, 0x2c, 0x20, 0x30, 0x78, 0x38, 0x34, 0x43, 0x38, 0x37, 0x38, 0x31, 0x34, 0x2c, 0x20, 0x30, 0x78, 0x38, 0x43, 0x43, 0x37, 0x30, 0x32, 0x30, 0x38, 0x2c, 0x20, 0x30, 0x78, 0x39, 0x30, 0x42, 0x45, 0x46, 0x46, 0x46, 0x41, 0x2c, 0x20, 0x30, 0x78, 0x41, 0x34, 0x35, 0x30, 0x36, 0x43, 0x45, 0x42, 0x2c, 0x20, 0x30, 0x78, 0x42, 0x45, 0x46, 0x39, 0x41, 0x33, 0x46, 0x37, 0x2c, 0x20, 0x30, 0x78, 0x43, 0x36, 0x37, 0x31, 0x37, 0x38, 0x46, 0x32, 0xd, 0xa, 0x9, 0x7d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x6d, 0x5b, 0x36, 0x34, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x31, 0x36, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x6d, 0x5b, 0x69, 0x5d, 0x20, 0x3d, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x31, 0x36, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x36, 0x34, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x6d, 0x5b, 0x69, 0x5d, 0x20, 0x3d, 0x20, 0x53, 0x49, 0x47, 0x31, 0x28, 0x6d, 0x5b, 0x69, 0x20, 0x2d, 0x20, 0x32, 0x5d, 0x29, 0x20, 0x2b, 0x20, 0x6d, 0x5b, 0x69, 0x20, 0x2d, 0x20, 0x37, 0x5d, 0x20, 0x2b, 0x20, 0x53, 0x49, 0x47, 0x30, 0x28, 0x6d, 0x5b, 0x69, 0x20, 0x2d, 0x20, 0x31, 0x35, 0x5d, 0x29, 0x20, 0x2b, 0x20, 0x6d, 0x5b, 0x69, 0x20, 0x2d, 0x20, 0x31, 0x36, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x61, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x30, 0x5d, 0x2c, 0x20, 0x62, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x31, 0x5d, 0x2c, 0x20, 0x63, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x32, 0x5d, 0x2c, 0x20, 0x64, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x33, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x65, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x34, 0x5d, 0x2c, 0x20, 0x66, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x35, 0x5d, 0x2c, 0x20, 0x67, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x36, 0x5d, 0x2c, 0x20, 0x68, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x37, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x43, 0x55, 0x44, 0x41, 0x5f, 0x41, 0x52, 0x43, 0x48, 0x5f, 0x5f, 0x29, 0x20, 0x7c, 0x7c, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4f, 0x50, 0x45, 0x4e, 0x43, 0x4c, 0x5f, 0x56, 0x45, 0x52, 0x53, 0x49, 0x4f, 0x4e, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x9, 0x23, 0x70, 0x72, 0x61, 0x67, 0x6d, 0x61, 0x20, 0x75, 0x6e, 0x72, 0x6f, 0x6c, 0x6c, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x36, 0x34, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x7b, 0xd, 0xa, 0x9, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x74, 0x31, 0x20, 0x3d, 0x20, 0x68, 0x20, 0x2b, 0x20, 0x45, 0x50, 0x31, 0x28, 0x65, 0x29, 0x20, 0x2b, 0x20, 0x43, 0x48, 0x28, 0x65, 0x2c, 0x20, 0x66, 0x2c, 0x20, 0x67, 0x29, 0x20, 0x2b, 0x20, 0x6b, 0x73, 0x68, 0x61, 0x5b, 0x69, 0x5d, 0x20, 0x2b, 0x20, 0x6d, 0x5b, 0x69, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x74, 0x32, 0x20, 0x3d, 0x20, 0x45, 0x50, 0x30, 0x28, 0x61, 0x29, 0x20, 0x2b, 0x20, 0x4d, 0x41, 0x4a, 0x28, 0x61, 0x2c, 0x20, 0x62, 0x2c, 0x20, 0x63, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x68, 0x20, 0x3d, 0x20, 0x67, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x67, 0x20, 0x3d, 0x20, 0x66, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x66, 0x20, 0x3d, 0x20, 0x65, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x65, 0x20, 0x3d, 0x20, 0x64, 0x20, 0x2b, 0x20, 0x74, 0x31, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x64, 0x20, 0x3d, 0x20, 0x63, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x63, 0x20, 0x3d, 0x20, 0x62, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x62, 0x20, 0x3d, 0x20, 0x61, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x61, 0x20, 0x3d, 0x20, 0x74, 0x31, 0x20, 0x2b, 0x20, 0x74, 0x32, 0x3b, 0xd, 0xa, 0x9, 0x7d, 0xd, 0xa, 0xd, 0xa, 0x9, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x30, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x61, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x31, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x62, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x32, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x63, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x33, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x64, 0x3b, 0xd, 0xa, 0x9, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x34, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x65, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x35, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x66, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x36, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x67, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x37, 0x5d, 0x20, 0x2b, 0x3d, 0x20, 0x68, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x21, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4f, 0x50, 0x45, 0x4e, 0x43, 0x4c, 0x5f, 0x56, 0x45, 0x52, 0x53, 0x49, 0x4f, 0x4e, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x73, 0x74, 0x61, 0x74, 0x69, 0x63, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x76, 0x65, 0x63, 0x38, 0x75, 0x20, 0x68, 0x32, 0x35, 0x36, 0x20, 0x3d, 0x20, 0x7b, 0xd, 0xa, 0x9, 0x30, 0x78, 0x36, 0x41, 0x30, 0x39, 0x45, 0x36, 0x36, 0x37, 0x2c, 0x20, 0x30, 0x78, 0x42, 0x42, 0x36, 0x37, 0x41, 0x45, 0x38, 0x35, 0x2c, 0x20, 0x30, 0x78, 0x33, 0x43, 0x36, 0x45, 0x46, 0x33, 0x37, 0x32, 0x2c, 0x20, 0x30, 0x78, 0x41, 0x35, 0x34, 0x46, 0x46, 0x35, 0x33, 0x41, 0x2c, 0x20, 0x30, 0x78, 0x35, 0x31, 0x30, 0x45, 0x35, 0x32, 0x37, 0x46, 0x2c, 0x20, 0x30, 0x78, 0x39, 0x42, 0x30, 0x35, 0x36, 0x38, 0x38, 0x43, 0x2c, 0x20, 0x30, 0x78, 0x31, 0x46, 0x38, 0x33, 0x44, 0x39, 0x41, 0x42, 0x2c, 0x20, 0x30, 0x78, 0x35, 0x42, 0x45, 0x30, 0x43, 0x44, 0x31, 0x39, 0xd, 0xa, 0x7d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x69, 0x6e, 0x6c, 0x69, 0x6e, 0x65, 0x20, 0x73, 0x74, 0x61, 0x74, 0x69, 0x63, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x2a, 0x70, 0x70, 0x2c, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x78, 0x29, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x38, 0x5f, 0x74, 0x20, 0x2a, 0x70, 0x20, 0x3d, 0x20, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x38, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x70, 0x70, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x70, 0x5b, 0x33, 0x5d, 0x20, 0x3d, 0x20, 0x78, 0x20, 0x26, 0x20, 0x30, 0x78, 0x66, 0x66, 0x3b, 0xd, 0xa, 0x9, 0x70, 0x5b, 0x32, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x78, 0x20, 0x3e, 0x3e, 0x20, 0x38, 0x29, 0x20, 0x26, 0x20, 0x30, 0x78, 0x66, 0x66, 0x3b, 0xd, 0xa, 0x9, 0x70, 0x5b, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x78, 0x20, 0x3e, 0x3e, 0x20, 0x31, 0x36, 0x29, 0x20, 0x26, 0x20, 0x30, 0x78, 0x66, 0x66, 0x3b, 0xd, 0xa, 0x9, 0x70, 0x5b, 0x30, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x78, 0x20, 0x3e, 0x3e, 0x20, 0x32, 0x34, 0x29, 0x20, 0x26, 0x20, 0x30, 0x78, 0x66, 0x66, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x69, 0x6e, 0x6c, 0x69, 0x6e, 0x65, 0x20, 0x73, 0x74, 0x61, 0x74, 0x69, 0x63, 0x20, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0x20, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x70, 0x72, 0x65, 0x70, 0x61, 0x72, 0x65, 0x5f, 0x6d, 0x73, 0x67, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x73, 0x2c, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x65, 0x78, 0x70, 0x69, 0x72, 0x65, 0x64, 0x2c, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x20, 0x63, 0x68, 0x61, 0x72, 0x20, 0x2a, 0x64, 0x61, 0x74, 0x61, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x9, 0x9, 0x9, 0x9, 0x9, 0x9, 0x9, 0x9, 0x9, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x2a, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x2c, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x20, 0x63, 0x68, 0x61, 0x72, 0x20, 0x2a, 0x72, 0x64, 0x61, 0x74, 0x61, 0x29, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x2f, 0x2a, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x73, 0x3a, 0x20, 0x25, 0x75, 0x2c, 0x20, 0x65, 0x78, 0x70, 0x69, 0x72, 0x65, 0x64, 0x3a, 0x20, 0x25, 0x30, 0x34, 0x78, 0x5c, 0x6e, 0x22, 0x2c, 0x20, 0x28, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x29, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x73, 0x2c, 0x20, 0x28, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x29, 0x65, 0x78, 0x70, 0x69, 0x72, 0x65, 0x64, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x64, 0x61, 0x74, 0x61, 0x3a, 0x5c, 0x6e, 0x22, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x20, 0x63, 0x68, 0x61, 0x72, 0x20, 0x70, 0x61, 0x64, 0x64, 0x65, 0x64, 0x5f, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x31, 0x32, 0x38, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x6d, 0x65, 0x6d, 0x73, 0x65, 0x74, 0x28, 0x70, 0x61, 0x64, 0x64, 0x65, 0x64, 0x5f, 0x64, 0x61, 0x74, 0x61, 0x2c, 0x20, 0x30, 0x2c, 0x20, 0x73, 0x69, 0x7a, 0x65, 0x6f, 0x66, 0x28, 0x70, 0x61, 0x64, 0x64, 0x65, 0x64, 0x5f, 0x64, 0x61, 0x74, 0x61, 0x29, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x6d, 0x65, 0x6d, 0x63, 0x70, 0x79, 0x28, 0x70, 0x61, 0x64, 0x64, 0x65, 0x64, 0x5f, 0x64, 0x61, 0x74, 0x61, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2c, 0x20, 0x31, 0x32, 0x33, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x33, 0x32, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x7b, 0xd, 0xa, 0x9, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x6a, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x6a, 0x20, 0x3c, 0x20, 0x34, 0x3b, 0x20, 0x6a, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x25, 0x30, 0x32, 0x68, 0x68, 0x78, 0x20, 0x22, 0x2c, 0x20, 0x70, 0x61, 0x64, 0x64, 0x65, 0x64, 0x5f, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x69, 0x20, 0x2a, 0x20, 0x34, 0x20, 0x2b, 0x20, 0x6a, 0x5d, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x5c, 0x6e, 0x22, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x7d, 0xd, 0xa, 0xd, 0xa, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x3a, 0x5c, 0x6e, 0x22, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x33, 0x32, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x25, 0x30, 0x32, 0x68, 0x68, 0x78, 0x20, 0x22, 0x2c, 0x20, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x20, 0x63, 0x68, 0x61, 0x72, 0x20, 0x2a, 0x29, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x29, 0x5b, 0x69, 0x5d, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x5c, 0x6e, 0x22, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x72, 0x64, 0x61, 0x74, 0x61, 0x3a, 0x5c, 0x6e, 0x22, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x73, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x7b, 0xd, 0xa, 0x9, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x6a, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x6a, 0x20, 0x3c, 0x20, 0x33, 0x32, 0x3b, 0x20, 0x6a, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x25, 0x30, 0x32, 0x68, 0x68, 0x78, 0x20, 0x22, 0x2c, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x69, 0x20, 0x2a, 0x20, 0x33, 0x32, 0x20, 0x2b, 0x20, 0x6a, 0x5d, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0x22, 0x5c, 0x6e, 0x22, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x7d, 0x2a, 0x2f, 0xd, 0xa, 0xd, 0xa, 0x9, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0x20, 0x6d, 0x73, 0x67, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x33, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x6d, 0x73, 0x67, 0x2e, 0x70, 0x73, 0x65, 0x65, 0x64, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x28, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x37, 0x36, 0x29, 0x20, 0x2b, 0x20, 0x69, 0x29, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x38, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x6d, 0x73, 0x67, 0x2e, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x20, 0x2b, 0x20, 0x69, 0x29, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x20, 0x3c, 0x20, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x73, 0x3b, 0x20, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x7b, 0xd, 0xa, 0x9, 0x9, 0x54, 0x68, 0x72, 0x65, 0x61, 0x64, 0x44, 0x61, 0x74, 0x61, 0x20, 0x2a, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3d, 0x20, 0x26, 0x6d, 0x73, 0x67, 0x2e, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x6e, 0x73, 0x69, 0x67, 0x6e, 0x65, 0x64, 0x20, 0x63, 0x68, 0x61, 0x72, 0x20, 0x2a, 0x74, 0x68, 0x72, 0x5f, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3d, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x67, 0x70, 0x75, 0x5f, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x20, 0x2a, 0x20, 0x33, 0x32, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x76, 0x65, 0x63, 0x38, 0x75, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x20, 0x3d, 0x20, 0x68, 0x32, 0x35, 0x36, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x76, 0x65, 0x63, 0x31, 0x36, 0x75, 0x20, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x30, 0x5d, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x30, 0x78, 0x30, 0x29, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x77, 0x6f, 0x72, 0x64, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x77, 0x6f, 0x72, 0x64, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x30, 0x78, 0x31, 0x29, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x77, 0x6f, 0x72, 0x64, 0x20, 0x26, 0x20, 0x30, 0x78, 0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x30, 0x30, 0x29, 0x20, 0x7c, 0x20, 0x28, 0x65, 0x78, 0x70, 0x69, 0x72, 0x65, 0x64, 0x20, 0x3e, 0x3e, 0x20, 0x32, 0x34, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x32, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x65, 0x78, 0x70, 0x69, 0x72, 0x65, 0x64, 0x20, 0x3c, 0x3c, 0x20, 0x38, 0x29, 0x20, 0x7c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x31, 0x31, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x33, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x3d, 0x20, 0x30, 0x78, 0x39, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x69, 0x29, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x77, 0x6f, 0x72, 0x64, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x30, 0x78, 0x41, 0x29, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x41, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x77, 0x6f, 0x72, 0x64, 0x20, 0x26, 0x20, 0x30, 0x78, 0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x30, 0x30, 0x29, 0x20, 0x7c, 0x20, 0x74, 0x68, 0x72, 0x5f, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x30, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x42, 0x2c, 0x20, 0x6a, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x3d, 0x20, 0x30, 0x78, 0x46, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x2c, 0x20, 0x6a, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x28, 0x74, 0x68, 0x72, 0x5f, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x6a, 0x29, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x73, 0x74, 0x61, 0x74, 0x65, 0x20, 0x3d, 0x20, 0x73, 0x68, 0x61, 0x32, 0x35, 0x36, 0x5f, 0x74, 0x72, 0x61, 0x6e, 0x73, 0x66, 0x6f, 0x72, 0x6d, 0x28, 0x73, 0x68, 0x61, 0x64, 0x61, 0x74, 0x61, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x2c, 0x20, 0x6a, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x3d, 0x20, 0x30, 0x78, 0x31, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x2c, 0x20, 0x6a, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x28, 0x74, 0x68, 0x72, 0x5f, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x32, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x6a, 0x29, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x77, 0x6f, 0x72, 0x64, 0x2c, 0x20, 0x2a, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x28, 0x74, 0x68, 0x72, 0x5f, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x32, 0x38, 0x29, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x30, 0x78, 0x32, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x77, 0x6f, 0x72, 0x64, 0x20, 0x3c, 0x3c, 0x20, 0x38, 0x29, 0x20, 0x7c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x37, 0x35, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x77, 0x6f, 0x72, 0x64, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x30, 0x78, 0x31, 0x36, 0x29, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x30, 0x78, 0x33, 0x5d, 0x20, 0x3d, 0x20, 0x28, 0x77, 0x6f, 0x72, 0x64, 0x20, 0x26, 0x20, 0x30, 0x78, 0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x30, 0x30, 0x29, 0x20, 0x7c, 0x20, 0x74, 0x68, 0x72, 0x5f, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x30, 0x5d, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x34, 0x2c, 0x20, 0x6a, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x3d, 0x20, 0x30, 0x78, 0x38, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x2c, 0x20, 0x6a, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x62, 0x65, 0x33, 0x32, 0x65, 0x6e, 0x63, 0x28, 0x26, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x69, 0x5d, 0x2c, 0x20, 0x2a, 0x28, 0x28, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x2a, 0x29, 0x28, 0x74, 0x68, 0x72, 0x5f, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x2b, 0x20, 0x31, 0x29, 0x20, 0x2b, 0x20, 0x6a, 0x29, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x7d, 0xd, 0xa, 0xd, 0xa, 0x9, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x20, 0x6d, 0x73, 0x67, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x74, 0x79, 0x70, 0x65, 0x64, 0x65, 0x66, 0x20, 0x73, 0x74, 0x72, 0x75, 0x63, 0x74, 0x20, 0x44, 0x65, 0x76, 0x48, 0x61, 0x73, 0x68, 0x52, 0x65, 0x73, 0x75, 0x6c, 0x74, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x20, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x3b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x76, 0x63, 0x70, 0x75, 0x2c, 0x20, 0x66, 0x6f, 0x75, 0x6e, 0x64, 0x3b, 0xd, 0xa, 0x7d, 0xd, 0xa, 0x44, 0x65, 0x76, 0x48, 0x61, 0x73, 0x68, 0x52, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x3b, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x43, 0x55, 0x44, 0x41, 0x43, 0x43, 0x5f, 0x5f, 0x29, 0x20, 0x7c, 0x7c, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4f, 0x50, 0x45, 0x4e, 0x43, 0x4c, 0x5f, 0x56, 0x45, 0x52, 0x53, 0x49, 0x4f, 0x4e, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4f, 0x50, 0x45, 0x4e, 0x43, 0x4c, 0x5f, 0x56, 0x45, 0x52, 0x53, 0x49, 0x4f, 0x4e, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x9, 0x5f, 0x5f, 0x6b, 0x65, 0x72, 0x6e, 0x65, 0x6c, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x67, 0x70, 0x75, 0x5f, 0x68, 0x61, 0x73, 0x68, 0x28, 0xd, 0xa, 0x9, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x20, 0x73, 0x74, 0x61, 0x72, 0x74, 0x5f, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x2c, 0x20, 0x5f, 0x5f, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0x20, 0x2a, 0x63, 0x5f, 0x6d, 0x73, 0x67, 0x2c, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x20, 0x44, 0x65, 0x76, 0x48, 0x61, 0x73, 0x68, 0x52, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x20, 0x2a, 0x72, 0x65, 0x73, 0x75, 0x6c, 0x74, 0xd, 0xa, 0x9, 0x29, 0xd, 0xa, 0x23, 0x65, 0x6c, 0x73, 0x65, 0xd, 0xa, 0x9, 0x5f, 0x5f, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x5f, 0x5f, 0x20, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0x20, 0x63, 0x5f, 0x6d, 0x73, 0x67, 0x3b, 0xd, 0xa, 0x9, 0x65, 0x78, 0x74, 0x65, 0x72, 0x6e, 0x20, 0x22, 0x43, 0x22, 0x20, 0x5f, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x5f, 0x20, 0x76, 0x6f, 0x69, 0x64, 0x20, 0x62, 0x69, 0x74, 0x63, 0x72, 0x65, 0x64, 0x69, 0x74, 0x5f, 0x67, 0x70, 0x75, 0x5f, 0x68, 0x61, 0x73, 0x68, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x20, 0x73, 0x74, 0x61, 0x72, 0x74, 0x5f, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x2c, 0x20, 0x44, 0x65, 0x76, 0x48, 0x61, 0x73, 0x68, 0x52, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x20, 0x2a, 0x72, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x29, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0x7b, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4f, 0x50, 0x45, 0x4e, 0x43, 0x4c, 0x5f, 0x56, 0x45, 0x52, 0x53, 0x49, 0x4f, 0x4e, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x20, 0x69, 0x64, 0x78, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x30, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x76, 0x63, 0x70, 0x75, 0x20, 0x3d, 0x20, 0x67, 0x65, 0x74, 0x5f, 0x67, 0x6c, 0x6f, 0x62, 0x61, 0x6c, 0x5f, 0x69, 0x64, 0x28, 0x31, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x5f, 0x5f, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0x20, 0x2a, 0x6d, 0x73, 0x67, 0x20, 0x3d, 0x20, 0x63, 0x5f, 0x6d, 0x73, 0x67, 0x3b, 0xd, 0xa, 0x9, 0x5f, 0x5f, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x61, 0x6e, 0x74, 0x20, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x54, 0x68, 0x72, 0x65, 0x61, 0x64, 0x44, 0x61, 0x74, 0x61, 0x20, 0x2a, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3d, 0x20, 0x26, 0x6d, 0x73, 0x67, 0x2d, 0x3e, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x76, 0x63, 0x70, 0x75, 0x5d, 0x3b, 0xd, 0xa, 0x23, 0x65, 0x6c, 0x73, 0x65, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x20, 0x69, 0x64, 0x78, 0x20, 0x3d, 0x20, 0x62, 0x6c, 0x6f, 0x63, 0x6b, 0x49, 0x64, 0x78, 0x2e, 0x78, 0x20, 0x2a, 0x20, 0x62, 0x6c, 0x6f, 0x63, 0x6b, 0x44, 0x69, 0x6d, 0x2e, 0x78, 0x20, 0x2b, 0x20, 0x74, 0x68, 0x72, 0x65, 0x61, 0x64, 0x49, 0x64, 0x78, 0x2e, 0x78, 0x3b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x76, 0x63, 0x70, 0x75, 0x20, 0x3d, 0x20, 0x62, 0x6c, 0x6f, 0x63, 0x6b, 0x49, 0x64, 0x78, 0x2e, 0x79, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x4d, 0x73, 0x67, 0x44, 0x61, 0x74, 0x61, 0x20, 0x2a, 0x6d, 0x73, 0x67, 0x20, 0x3d, 0x20, 0x26, 0x63, 0x5f, 0x6d, 0x73, 0x67, 0x3b, 0xd, 0xa, 0x9, 0x63, 0x6f, 0x6e, 0x73, 0x74, 0x20, 0x54, 0x68, 0x72, 0x65, 0x61, 0x64, 0x44, 0x61, 0x74, 0x61, 0x20, 0x2a, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3d, 0x20, 0x26, 0x63, 0x5f, 0x6d, 0x73, 0x67, 0x2e, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x76, 0x63, 0x70, 0x75, 0x5d, 0x3b, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0xd, 0xa, 0x9, 0x76, 0x65, 0x63, 0x38, 0x75, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x73, 0x74, 0x61, 0x74, 0x65, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x76, 0x65, 0x63, 0x31, 0x36, 0x75, 0x20, 0x64, 0x61, 0x74, 0x61, 0x3b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x20, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x20, 0x3d, 0x20, 0x73, 0x74, 0x61, 0x72, 0x74, 0x5f, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x20, 0x2b, 0x20, 0x69, 0x64, 0x78, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x36, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x30, 0x5d, 0x2c, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x37, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x31, 0x5d, 0x2c, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x38, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x32, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3d, 0x20, 0x28, 0x28, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x29, 0x72, 0x64, 0x61, 0x74, 0x61, 0x36, 0x20, 0x3c, 0x3c, 0x20, 0x35, 0x36, 0x29, 0x20, 0x7c, 0x20, 0x28, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x29, 0x72, 0x64, 0x61, 0x74, 0x61, 0x37, 0x20, 0x3c, 0x3c, 0x20, 0x32, 0x34, 0x29, 0x20, 0x7c, 0x20, 0x28, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x36, 0x34, 0x5f, 0x74, 0x29, 0x72, 0x64, 0x61, 0x74, 0x61, 0x38, 0x20, 0x3e, 0x3e, 0x20, 0x38, 0x29, 0x29, 0x20, 0x2b, 0x20, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x72, 0x64, 0x61, 0x74, 0x61, 0x36, 0x20, 0x3d, 0x20, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x29, 0x28, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3e, 0x3e, 0x20, 0x35, 0x36, 0x29, 0x20, 0x7c, 0x20, 0x28, 0x72, 0x64, 0x61, 0x74, 0x61, 0x36, 0x20, 0x26, 0x20, 0x30, 0x78, 0x46, 0x46, 0x46, 0x46, 0x46, 0x46, 0x30, 0x30, 0x29, 0x3b, 0xd, 0xa, 0x9, 0x72, 0x64, 0x61, 0x74, 0x61, 0x37, 0x20, 0x3d, 0x20, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x29, 0x28, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3e, 0x3e, 0x20, 0x32, 0x34, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x31, 0x30, 0x20, 0x3d, 0x20, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x29, 0x28, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3c, 0x3c, 0x20, 0x38, 0x29, 0x20, 0x7c, 0x20, 0x30, 0x78, 0x38, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x72, 0x64, 0x61, 0x74, 0x61, 0x38, 0x20, 0x3d, 0x20, 0x28, 0x75, 0x69, 0x6e, 0x74, 0x33, 0x32, 0x5f, 0x74, 0x29, 0x28, 0x72, 0x64, 0x61, 0x74, 0x61, 0x20, 0x3c, 0x3c, 0x20, 0x38, 0x29, 0x20, 0x7c, 0x20, 0x28, 0x72, 0x64, 0x61, 0x74, 0x61, 0x38, 0x20, 0x26, 0x20, 0x30, 0x78, 0x46, 0x46, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x30, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x36, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x37, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x32, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x38, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x33, 0x5d, 0x20, 0x3d, 0x20, 0x6d, 0x73, 0x67, 0x2d, 0x3e, 0x70, 0x73, 0x65, 0x65, 0x64, 0x5b, 0x30, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x34, 0x5d, 0x20, 0x3d, 0x20, 0x6d, 0x73, 0x67, 0x2d, 0x3e, 0x70, 0x73, 0x65, 0x65, 0x64, 0x5b, 0x31, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x35, 0x5d, 0x20, 0x3d, 0x20, 0x6d, 0x73, 0x67, 0x2d, 0x3e, 0x70, 0x73, 0x65, 0x65, 0x64, 0x5b, 0x32, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x36, 0x5d, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x33, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x37, 0x5d, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x34, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x38, 0x5d, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x35, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x39, 0x5d, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x36, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x41, 0x5d, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x37, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x42, 0x5d, 0x20, 0x3d, 0x20, 0x74, 0x68, 0x72, 0x64, 0x61, 0x74, 0x61, 0x2d, 0x3e, 0x72, 0x64, 0x61, 0x74, 0x61, 0x5b, 0x38, 0x5d, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x43, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x36, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x44, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x37, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x45, 0x5d, 0x20, 0x3d, 0x20, 0x72, 0x64, 0x61, 0x74, 0x61, 0x31, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x46, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x2f, 0x2a, 0x69, 0x66, 0x20, 0x28, 0x76, 0x63, 0x70, 0x75, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x20, 0x26, 0x26, 0x20, 0x69, 0x64, 0x78, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x29, 0xd, 0xa, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0xd, 0xa, 0x9, 0x9, 0x9, 0x22, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x5c, 0x6e, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x5c, 0x6e, 0x5c, 0x6e, 0x22, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x32, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x33, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x34, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x35, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x36, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x37, 0x5d, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x38, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x39, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x30, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x31, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x32, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x33, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x34, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x35, 0x5d, 0xd, 0xa, 0x9, 0x9, 0x29, 0x3b, 0x2a, 0x2f, 0xd, 0xa, 0xd, 0xa, 0x9, 0x73, 0x74, 0x61, 0x74, 0x65, 0x20, 0x3d, 0x20, 0x73, 0x68, 0x61, 0x32, 0x35, 0x36, 0x5f, 0x74, 0x72, 0x61, 0x6e, 0x73, 0x66, 0x6f, 0x72, 0x6d, 0x28, 0x64, 0x61, 0x74, 0x61, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x30, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x31, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x32, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x33, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x34, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x35, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x36, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x37, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x38, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x39, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x41, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x42, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x43, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x44, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x45, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x30, 0x3b, 0xd, 0xa, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x78, 0x46, 0x5d, 0x20, 0x3d, 0x20, 0x30, 0x78, 0x30, 0x30, 0x30, 0x30, 0x30, 0x33, 0x64, 0x38, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x2f, 0x2a, 0x69, 0x66, 0x20, 0x28, 0x76, 0x63, 0x70, 0x75, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x20, 0x26, 0x26, 0x20, 0x69, 0x64, 0x78, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x29, 0xd, 0xa, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0xd, 0xa, 0x9, 0x9, 0x9, 0x22, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x5c, 0x6e, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x5c, 0x6e, 0x5c, 0x6e, 0x22, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x30, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x32, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x33, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x34, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x35, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x36, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x37, 0x5d, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x9, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x38, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x39, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x30, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x31, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x32, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x33, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x34, 0x5d, 0x2c, 0x20, 0x64, 0x61, 0x74, 0x61, 0x2e, 0x76, 0x5b, 0x31, 0x35, 0x5d, 0xd, 0xa, 0x9, 0x9, 0x29, 0x3b, 0x2a, 0x2f, 0xd, 0xa, 0xd, 0xa, 0x9, 0x73, 0x74, 0x61, 0x74, 0x65, 0x20, 0x3d, 0x20, 0x73, 0x68, 0x61, 0x32, 0x35, 0x36, 0x5f, 0x74, 0x72, 0x61, 0x6e, 0x73, 0x66, 0x6f, 0x72, 0x6d, 0x28, 0x64, 0x61, 0x74, 0x61, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x29, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x2f, 0x2a, 0x69, 0x66, 0x20, 0x28, 0x76, 0x63, 0x70, 0x75, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x20, 0x26, 0x26, 0x20, 0x69, 0x64, 0x78, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x29, 0xd, 0xa, 0x9, 0x9, 0x70, 0x72, 0x69, 0x6e, 0x74, 0x66, 0x28, 0xd, 0xa, 0x9, 0x9, 0x9, 0x22, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x20, 0x25, 0x30, 0x34, 0x78, 0x5c, 0x6e, 0x5c, 0x6e, 0x22, 0x2c, 0xd, 0xa, 0x9, 0x9, 0x9, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x30, 0x5d, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x31, 0x5d, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x32, 0x5d, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x33, 0x5d, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x34, 0x5d, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x35, 0x5d, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x36, 0x5d, 0x2c, 0x20, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x37, 0x5d, 0xd, 0xa, 0x9, 0x9, 0x29, 0x3b, 0x2a, 0x2f, 0xd, 0xa, 0xd, 0xa, 0x9, 0x66, 0x6f, 0x72, 0x20, 0x28, 0x69, 0x6e, 0x74, 0x20, 0x69, 0x20, 0x3d, 0x20, 0x30, 0x3b, 0x20, 0x69, 0x20, 0x3c, 0x20, 0x38, 0x3b, 0x20, 0x69, 0x20, 0x2b, 0x3d, 0x20, 0x31, 0x29, 0xd, 0xa, 0x9, 0x7b, 0xd, 0xa, 0x9, 0x9, 0x69, 0x66, 0x20, 0x28, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x20, 0x3e, 0x20, 0x6d, 0x73, 0x67, 0x2d, 0x3e, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x69, 0x66, 0x20, 0x28, 0x73, 0x74, 0x61, 0x74, 0x65, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x20, 0x3c, 0x20, 0x6d, 0x73, 0x67, 0x2d, 0x3e, 0x74, 0x61, 0x72, 0x67, 0x65, 0x74, 0x2e, 0x76, 0x5b, 0x69, 0x5d, 0x29, 0xd, 0xa, 0x9, 0x9, 0x7b, 0xd, 0xa, 0x23, 0x69, 0x66, 0x20, 0x64, 0x65, 0x66, 0x69, 0x6e, 0x65, 0x64, 0x28, 0x5f, 0x5f, 0x4f, 0x50, 0x45, 0x4e, 0x43, 0x4c, 0x5f, 0x56, 0x45, 0x52, 0x53, 0x49, 0x4f, 0x4e, 0x5f, 0x5f, 0x29, 0xd, 0xa, 0x9, 0x9, 0x9, 0x69, 0x66, 0x20, 0x28, 0x61, 0x74, 0x6f, 0x6d, 0x69, 0x63, 0x5f, 0x61, 0x64, 0x64, 0x28, 0x26, 0x72, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x2d, 0x3e, 0x66, 0x6f, 0x75, 0x6e, 0x64, 0x2c, 0x20, 0x31, 0x29, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x29, 0xd, 0xa, 0x23, 0x65, 0x6c, 0x73, 0x65, 0xd, 0xa, 0x9, 0x9, 0x9, 0x69, 0x66, 0x20, 0x28, 0x61, 0x74, 0x6f, 0x6d, 0x69, 0x63, 0x41, 0x64, 0x64, 0x28, 0x26, 0x72, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x2d, 0x3e, 0x66, 0x6f, 0x75, 0x6e, 0x64, 0x2c, 0x20, 0x31, 0x29, 0x20, 0x3d, 0x3d, 0x20, 0x30, 0x29, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0x9, 0x9, 0x9, 0x9, 0x72, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x2d, 0x3e, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x20, 0x3d, 0x20, 0x6e, 0x6f, 0x6e, 0x63, 0x65, 0x2c, 0x20, 0x72, 0x65, 0x73, 0x75, 0x6c, 0x74, 0x2d, 0x3e, 0x76, 0x63, 0x70, 0x75, 0x20, 0x3d, 0x20, 0x76, 0x63, 0x70, 0x75, 0x3b, 0xd, 0xa, 0xd, 0xa, 0x9, 0x9, 0x9, 0x72, 0x65, 0x74, 0x75, 0x72, 0x6e, 0x3b, 0xd, 0xa, 0x9, 0x9, 0x7d, 0xd, 0xa, 0x9, 0x7d, 0xd, 0xa, 0x7d, 0xd, 0xa, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa, 0xd, 0xa, 0xd, 0xa, 0x23, 0x65, 0x6e, 0x64, 0x69, 0x66, 0xd, 0xa}; +unsigned int cruncher_h_len = 9417; diff --git a/crypto/util/opencl/opencl.cpp b/crypto/util/opencl/opencl.cpp index a401e599..69a6d478 100644 --- a/crypto/util/opencl/opencl.cpp +++ b/crypto/util/opencl/opencl.cpp @@ -2,6 +2,7 @@ #include "opencl.h" #include "miner.h" #include "opencl_helper.h" +#include "../cruncher.h" #ifdef __APPLE__ #include @@ -12,16 +13,17 @@ namespace opencl { void OpenCL::load_source(const char *filename) { - FILE *fp; + FILE *fp = fopen(filename, "r"); - fp = fopen(filename, "r"); if (!fp) { LOG(ERROR) << "[ OpenCL: failed to load kernel source '" << filename << "' ]"; exit(1); } + source_str_ = (char *)malloc(MAX_SOURCE_SIZE); source_size_ = fread(source_str_, 1, MAX_SOURCE_SIZE, fp); fclose(fp); + if (GET_VERBOSITY_LEVEL() >= VERBOSITY_NAME(INFO)) { LOG(PLAIN) << "[ OpenCL: loaded kernel source '" << filename << "' (" << source_size_ << " bytes) ]"; } @@ -31,6 +33,7 @@ void OpenCL::set_source(unsigned char *source, unsigned int length) { source_str_ = (char *)malloc(MAX_SOURCE_SIZE); memcpy(source_str_, source, length); source_size_ = length; + if (GET_VERBOSITY_LEVEL() >= VERBOSITY_NAME(INFO)) { LOG(PLAIN) << "[ OpenCL: set kernel source (" << source_size_ << " bytes) ]"; } @@ -47,6 +50,7 @@ void OpenCL::print_devices() { // devices char buf[1024]; num_devices_ = 0; + for (uint p = 0; p < platform_count_; p++) { cl_err = clGetDeviceIDs(platforms_[p], CL_DEVICE_TYPE_ALL, 0, NULL, &device_count_); if (cl_err != CL_SUCCESS) { @@ -57,6 +61,7 @@ void OpenCL::print_devices() { devices_ = (cl_device_id *)malloc(device_count_ * sizeof(cl_device_id)); CL_WRAPPER(clGetDeviceIDs(platforms_[p], CL_DEVICE_TYPE_ALL, device_count_, devices_, NULL)); + for (uint i = 0; i < device_count_; i++) { CL_WRAPPER(clGetDeviceInfo(devices_[i], CL_DEVICE_NAME, sizeof(buf), buf, NULL)); if (GET_VERBOSITY_LEVEL() >= VERBOSITY_NAME(INFO)) { @@ -78,11 +83,11 @@ void OpenCL::create_context(cl_uint platform_idx, cl_uint device_idx) { free(devices_); CL_WRAPPER(clGetDeviceIDs(platforms_[platform_idx], CL_DEVICE_TYPE_ALL, 0, NULL, &device_count_)); - devices_ = (cl_device_id *)malloc(device_count_ * sizeof(cl_device_id)); CL_WRAPPER(clGetDeviceIDs(platforms_[platform_idx], CL_DEVICE_TYPE_ALL, device_count_, devices_, NULL)); CL_WRAPPER(clGetDeviceInfo(devices_[device_idx], CL_DEVICE_NAME, sizeof(buf), buf, NULL)); + CL_WRAPPER(clGetDeviceInfo(devices_[device_idx], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_work_group_size_), &max_work_group_size_, NULL)); @@ -100,91 +105,50 @@ void OpenCL::create_context(cl_uint platform_idx, cl_uint device_idx) { } void OpenCL::create_kernel() { - // printf("[ OpenCL: create kernel ]\n"); + // printf("[ OpenCL: create kernel ]\n"); cl_int ret; program_ = clCreateProgramWithSource(context_, 1, (const char **)&source_str_, (const size_t *)&source_size_, &ret); CL_WRAPPER(ret); + ret = clBuildProgram(program_, 1, &devices_[device_idx_], NULL, NULL, NULL); if (ret != CL_SUCCESS) { size_t blen = 0; CL_WRAPPER(clGetProgramBuildInfo(program_, devices_[device_idx_], CL_PROGRAM_BUILD_LOG, 0, NULL, &blen)); - char *buffer = (char *) malloc(blen); - - CL_WRAPPER(clGetProgramBuildInfo(program_, devices_[device_idx_], CL_PROGRAM_BUILD_LOG, blen, &buffer, NULL)); + char *buffer = (char *)malloc(blen + 1); + buffer[blen] = '\0'; + CL_WRAPPER(clGetProgramBuildInfo(program_, devices_[device_idx_], CL_PROGRAM_BUILD_LOG, blen, buffer, NULL)); LOG(ERROR) << "[ OpenCL: ERROR ]\n" << buffer << "\n"; + free(buffer); exit(4); } - kernel_ = clCreateKernel(program_, "sha256", &ret); + kernel_ = clCreateKernel(program_, "bitcredit_gpu_hash", &ret); CL_WRAPPER(ret); + command_queue_ = clCreateCommandQueue(context_, devices_[device_idx_], 0, &ret); CL_WRAPPER(ret); } -void OpenCL::load_objects(uint32_t gpu_id, uint32_t cpu_id, unsigned char *data, const uint8_t *target, - unsigned char *rdata, uint32_t gpu_threads) { - - static const int len = 123, n = 3; - - uint32_t PaddedMessage[16 * n]; // bring balance to the force, 512*3 bits - memset(PaddedMessage, 0, 16 * n * sizeof(uint32_t)); - memcpy(PaddedMessage, data, len); - ((uchar *)PaddedMessage)[len] = 0x80; // guard bit after data - uint32_t endiandata[16 * n]; - for (int k = 0; k < 16 * n; k++) - be32enc(&endiandata[k], ((uint32_t *)PaddedMessage)[k]); - ((uint32_t *)endiandata)[16 * n - 1] = len * 8; // size to the end - - uint32_t endiantarget[8]; - for (int k = 0; k < 8; k++) - be32enc(&endiantarget[k], ((uint32_t *)target)[k]); - - // std::cout << "PaddedMessage[" << 16 * n * sizeof(uint32_t) << "]: "; - // for (int z = 0; z < 16 * n; z++) - // printf("%08x ", endiandata[z]); - // std::cout << std::endl; +void OpenCL::load_objects(uint32_t gpu_id, uint32_t cpu_id, uint32_t expired, const unsigned char *data, + const uint8_t *target, const unsigned char *rdata, uint32_t gpu_threads) { + (void)cpu_id, (void)gpu_id; + MsgData msg = bitcredit_prepare_msg(gpu_threads, expired, data, target, rdata); cl_int ret; - buffer_rdata_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_uchar) * 32 * gpu_threads, NULL, &ret); - CL_WRAPPER(ret); - buffer_data_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_uint) * 16 * 3, NULL, &ret); - CL_WRAPPER(ret); - buffer_target_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_uint) * 8, NULL, &ret); - CL_WRAPPER(ret); - buffer_gpu_threads_= clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_uint), NULL, &ret); - CL_WRAPPER(ret); - buffer_cpu_id_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_uint), NULL, &ret); - CL_WRAPPER(ret); - buffer_threads_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_ulong), NULL, &ret); - CL_WRAPPER(ret); - buffer_start_nonce_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_ulong), NULL, &ret); - CL_WRAPPER(ret); - buffer_expired_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(cl_uint), NULL, &ret); + + buffer_msg_ = clCreateBuffer(context_, CL_MEM_READ_ONLY, sizeof(msg), NULL, &ret); CL_WRAPPER(ret); - buffer_result_ = clCreateBuffer(context_, CL_MEM_READ_WRITE, sizeof(cl_ulong) * 2, NULL, &ret); + + buffer_result_ = clCreateBuffer(context_, CL_MEM_READ_WRITE, sizeof(DevHashResult), NULL, &ret); CL_WRAPPER(ret); - CL_WRAPPER(clEnqueueWriteBuffer(command_queue_, buffer_rdata_, CL_TRUE, 0, sizeof(cl_uchar) * 32 * gpu_threads, rdata, 0, - NULL, NULL)); - CL_WRAPPER(clEnqueueWriteBuffer(command_queue_, buffer_data_, CL_TRUE, 0, sizeof(cl_uint) * 16 * 3, endiandata, 0, - NULL, NULL)); - CL_WRAPPER(clEnqueueWriteBuffer(command_queue_, buffer_target_, CL_TRUE, 0, sizeof(cl_uint) * 8, endiantarget, 0, - NULL, NULL)); - CL_WRAPPER(clEnqueueWriteBuffer(command_queue_, buffer_gpu_threads_, CL_TRUE, 0, sizeof(cl_uint), &gpu_threads, 0, - NULL, NULL)); - - CL_WRAPPER(clSetKernelArg(kernel_, 0, sizeof(buffer_rdata_), (void *)&buffer_rdata_)); - CL_WRAPPER(clSetKernelArg(kernel_, 1, sizeof(buffer_data_), (void *)&buffer_data_)); - CL_WRAPPER(clSetKernelArg(kernel_, 2, sizeof(buffer_target_), (void *)&buffer_target_)); - CL_WRAPPER(clSetKernelArg(kernel_, 3, sizeof(buffer_gpu_threads_), (void *)&buffer_gpu_threads_)); - CL_WRAPPER(clSetKernelArg(kernel_, 4, sizeof(buffer_cpu_id_), (void *)&buffer_cpu_id_)); - CL_WRAPPER(clSetKernelArg(kernel_, 5, sizeof(buffer_threads_), (void *)&buffer_threads_)); - CL_WRAPPER(clSetKernelArg(kernel_, 6, sizeof(buffer_start_nonce_), (void *)&buffer_start_nonce_)); - CL_WRAPPER(clSetKernelArg(kernel_, 7, sizeof(buffer_expired_), (void *)&buffer_expired_)); - CL_WRAPPER(clSetKernelArg(kernel_, 8, sizeof(buffer_result_), (void *)&buffer_result_)); + CL_WRAPPER(clSetKernelArg(kernel_, 1, sizeof(buffer_msg_), (void *)&buffer_msg_)); + CL_WRAPPER(clSetKernelArg(kernel_, 2, sizeof(buffer_result_), (void *)&buffer_result_)); + + CL_WRAPPER(clEnqueueWriteBuffer(command_queue_, buffer_msg_, CL_TRUE, 0, sizeof(msg), &msg, 0, NULL, NULL)); } void OpenCL::release() { @@ -195,43 +159,36 @@ void OpenCL::release() { CL_WRAPPER(clReleaseDevice(devices_[device_idx_])); CL_WRAPPER(clReleaseMemObject(buffer_result_)); - CL_WRAPPER(clReleaseMemObject(buffer_expired_)); - CL_WRAPPER(clReleaseMemObject(buffer_start_nonce_)); - CL_WRAPPER(clReleaseMemObject(buffer_threads_)); - CL_WRAPPER(clReleaseMemObject(buffer_cpu_id_)); - CL_WRAPPER(clReleaseMemObject(buffer_gpu_threads_)); - CL_WRAPPER(clReleaseMemObject(buffer_target_)); - CL_WRAPPER(clReleaseMemObject(buffer_data_)); - CL_WRAPPER(clReleaseMemObject(buffer_rdata_)); + CL_WRAPPER(clReleaseMemObject(buffer_msg_)); free(devices_); free(platforms_); free(source_str_); } -HashResult OpenCL::scan_hash(uint cpu_id, uint32_t gpu_threads, td::uint64 threads, td::uint64 start_nonce, uint expired) { - td::uint64 start = (start_nonce / gpu_threads); - cl_ulong result[2] = {UINT64_MAX, UINT64_MAX}; +HashResult OpenCL::scan_hash(uint cpu_id, uint32_t gpu_threads, uint64_t threads, uint64_t start_nonce, uint expired) { + (void)expired; - CL_WRAPPER(clEnqueueWriteBuffer(command_queue_, buffer_cpu_id_, CL_TRUE, 0, sizeof(cl_uint), &cpu_id, 0, NULL, NULL)); - CL_WRAPPER( - clEnqueueWriteBuffer(command_queue_, buffer_threads_, CL_TRUE, 0, sizeof(cl_ulong), &threads, 0, NULL, NULL)); - CL_WRAPPER( - clEnqueueWriteBuffer(command_queue_, buffer_start_nonce_, CL_TRUE, 0, sizeof(cl_ulong), &start, 0, NULL, NULL)); - CL_WRAPPER( - clEnqueueWriteBuffer(command_queue_, buffer_expired_, CL_TRUE, 0, sizeof(cl_uint), &expired, 0, NULL, NULL)); - CL_WRAPPER(clEnqueueWriteBuffer(command_queue_, buffer_result_, CL_TRUE, 0, sizeof(result), result, 0, NULL, NULL)); + HashResult r; + r.nonce = UINT64_MAX, r.vcpu = UINT64_MAX, r.cpu_id = cpu_id; - size_t global_work_size[2] = {threads / gpu_threads, gpu_threads}; + uint64_t start = start_nonce / gpu_threads; + CL_WRAPPER(clSetKernelArg(kernel_, 0, sizeof(start), (void *)&start)); + + DevHashResult devresult; + devresult.nonce = UINT64_MAX, devresult.vcpu = UINT32_MAX, devresult.found = 0; + + clEnqueueWriteBuffer(command_queue_, buffer_result_, CL_FALSE, 0, sizeof(devresult), &devresult, 0, NULL, NULL); + size_t global_work_size[2] = {threads / gpu_threads, gpu_threads}; CL_WRAPPER(clEnqueueNDRangeKernel(command_queue_, kernel_, 2, NULL, global_work_size, NULL, 0, NULL, NULL)); + CL_WRAPPER(clFinish(command_queue_)); - clEnqueueReadBuffer(command_queue_, buffer_result_, CL_TRUE, 0, sizeof(result), result, 0, NULL, NULL); + clEnqueueReadBuffer(command_queue_, buffer_result_, CL_TRUE, 0, sizeof(devresult), &devresult, 0, NULL, NULL); + + r.nonce = devresult.nonce; + r.vcpu = (devresult.vcpu == UINT32_MAX) ? UINT64_MAX : devresult.vcpu; - HashResult r; - r.nonce = result[0]; - r.vcpu = result[1]; - r.cpu_id = cpu_id; return r; } diff --git a/crypto/util/opencl/opencl.h b/crypto/util/opencl/opencl.h index ea535196..f619c4ca 100644 --- a/crypto/util/opencl/opencl.h +++ b/crypto/util/opencl/opencl.h @@ -1,5 +1,6 @@ #include #include +#include #include #include @@ -16,6 +17,7 @@ typedef unsigned int uint; #endif namespace opencl { + struct HashResult { uint64_t nonce; uint64_t vcpu; @@ -31,8 +33,9 @@ class OpenCL { int get_num_devices(); void create_context(cl_uint platform_idx, cl_uint device_idx); void create_kernel(); - void load_objects(uint32_t gpu_id, uint32_t cpu_id, unsigned char *data, const uint8_t *target, unsigned char *rdata, uint32_t gpu_threads); - HashResult scan_hash(uint cpu_id, uint32_t gpu_threads, td::uint64 threads, td::uint64 start_nonce, uint expired); + void load_objects(uint32_t gpu_id, uint32_t cpu_id, uint32_t expired, const unsigned char *data, + const uint8_t *target, const unsigned char *rdata, uint32_t gpu_threads); + HashResult scan_hash(uint cpu_id, uint32_t gpu_threads, uint64_t threads, uint64_t start_nonce, uint expired); void release(); private: @@ -41,27 +44,25 @@ class OpenCL { private: size_t source_size_; char *source_str_; + cl_uint device_count_; cl_uint platform_count_; size_t max_work_group_size_; + cl_platform_id *platforms_; cl_device_id *devices_; + cl_uint platform_idx_; cl_uint device_idx_; + cl_context context_; cl_program program_; cl_kernel kernel_; cl_command_queue command_queue_; private: - cl_mem buffer_rdata_; - cl_mem buffer_data_; - cl_mem buffer_target_; - cl_mem buffer_gpu_threads_; - cl_mem buffer_cpu_id_; - cl_mem buffer_threads_; - cl_mem buffer_start_nonce_; - cl_mem buffer_expired_; + cl_mem buffer_msg_; cl_mem buffer_result_; }; + } // namespace opencl diff --git a/crypto/util/opencl/sha256.cpp b/crypto/util/opencl/sha256.cpp index c344ae75..107232f4 100644 --- a/crypto/util/opencl/sha256.cpp +++ b/crypto/util/opencl/sha256.cpp @@ -3,7 +3,7 @@ #include "sha256.h" #include "miner.h" #include "opencl.h" -#include "sha256_cl.h" +#include "../cruncher_h.h" #include "crypto/util/Miner.h" @@ -13,8 +13,8 @@ td::optional SHA256::run(ton::HDataEnv H, unsigned char *rdata, con int cpu_id) { // opencl auto opencl = OpenCL(); - //opencl.load_source("sha256.cl"); - opencl.set_source(sha256_cl, sha256_cl_len); + //opencl.load_source("cruncher.h"); + opencl.set_source(cruncher_h, cruncher_h_len); opencl.print_devices(); opencl.create_context(options.platform_id, options.gpu_id); opencl.create_kernel(); @@ -22,24 +22,25 @@ td::optional SHA256::run(ton::HDataEnv H, unsigned char *rdata, con // data td::Slice data = H.as_slice(); - td::uint64 throughput = (td::uint64)((1U << 19) * options.factor);// 256*256*64*8*factor/64 + td::uint64 throughput = (td::uint64)((1U << 19) * options.factor); // 256*256*64*8*factor/64 if (options.max_iterations < throughput) { throughput = options.max_iterations; } LOG(WARNING) << "[ START MINER, GPU ID: " << options.gpu_id << ", boost factor: " << options.factor << ", throughput: " << throughput << " ]"; - // set data + // set once at start + uint32_t expired = options.expire_base; + + // set data //std::cout << "data: " << hex_encode(data) << std::endl; unsigned char input[123], complexity[32]; memcpy(input, data.ubegin(), data.size()); - opencl.load_objects(options.gpu_id, cpu_id, input, options.complexity.data(), rdata, options.gpu_threads); + opencl.load_objects(options.gpu_id, cpu_id, expired, input, options.complexity.data(), rdata, options.gpu_threads); if (options.instant_hashes_computed) { *options.instant_hashes_computed = throughput; } - // set once at start - uint32_t expired = options.expire_base; td::int64 i = 0; for (; i < options.max_iterations;) { td::Timestamp instant_start_at = td::Timestamp::now();