From f3e8ffe2e99b3bf47a48bb22dd5e624d1701929e Mon Sep 17 00:00:00 2001 From: Xuanyi Zhou Date: Thu, 17 Sep 2020 23:47:26 -0400 Subject: [PATCH 1/4] all basic implementations --- stream_compaction/cpu.cu | 44 ++++++++++-- stream_compaction/efficient.cu | 123 +++++++++++++++++++++++++++++++-- stream_compaction/naive.cu | 41 +++++++++-- stream_compaction/thrust.cu | 7 +- 4 files changed, 200 insertions(+), 15 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..40b95f3 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -19,7 +19,13 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + + int total = 0; + for (int i = 0; i < n; ++i) { + odata[i] = total; + total += idata[i]; + } + timer().endCpuTimer(); } @@ -30,9 +36,17 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + + int count = 0; + for (int i = 0; i < n; ++i) { + if (idata[i] != 0) { + odata[count] = idata[i]; + ++count; + } + } + timer().endCpuTimer(); - return -1; + return count; } /** @@ -42,9 +56,29 @@ namespace StreamCompaction { */ int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + + // scan + int total = 0; + for (int i = 0; i < n; ++i) { + odata[i] = total; + total += (idata[i] == 0 ? 0 : 1); + } + + // scatter + int count = 0; + for (int i = 1; i < n; ++i) { + if (odata[i] != count) { + odata[count] = idata[i - 1]; + ++count; + } + } + if (idata[n - 1] != 0) { + odata[count] = idata[n - 1]; + ++count; + } + timer().endCpuTimer(); - return -1; + return count; } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..2ab90b5 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -6,19 +6,110 @@ namespace StreamCompaction { namespace Efficient { using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { + PerformanceTimer& timer() { static PerformanceTimer timer; return timer; } + __global__ void kernScanPerBlock(int *data, int *lastData) { + extern __shared__ int buffer[]; + data += blockIdx.x * blockDim.x; + + // copy data to shared memory + buffer[threadIdx.x] = data[threadIdx.x]; + __syncthreads(); + + int lastElem = 0; + if (lastData && threadIdx.x == blockDim.x - 1) { + lastElem = buffer[threadIdx.x]; + } + + // upward pass + for (int gap = 2; gap < blockDim.x; gap <<= 1) { + if ((threadIdx.x & (gap - 1)) == gap - 1) { + buffer[threadIdx.x] += buffer[threadIdx.x - (gap >> 1)]; + } + __syncthreads(); + } + + if (threadIdx.x == blockDim.x - 1) { + int halfIdx = threadIdx.x >> 1; + buffer[threadIdx.x] = buffer[halfIdx]; + buffer[halfIdx] = 0; + } + __syncthreads(); + + // downward pass + for (int gap = blockDim.x >> 1; gap > 1; gap >>= 1) { + if ((threadIdx.x & (gap - 1)) == gap - 1) { + int prevIdx = threadIdx.x - (gap >> 1); + int sum = buffer[threadIdx.x] + buffer[prevIdx]; + buffer[prevIdx] = buffer[threadIdx.x]; + buffer[threadIdx.x] = sum; + } + __syncthreads(); + } + + // copy data back + data[threadIdx.x] = buffer[threadIdx.x]; + if (lastData && threadIdx.x == blockDim.x - 1) { + lastData[blockIdx.x] = lastElem + buffer[threadIdx.x]; + } + } + + void _computeSizes(int n, int log2BlockSize, int *blockSize, int *numBlocks, int *bufferSize) { + *blockSize = 1 << log2BlockSize; + *numBlocks = n >> log2BlockSize; + if ((n & (*blockSize - 1)) != 0) { + ++*numBlocks; + } + *bufferSize = *numBlocks << log2BlockSize; + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + int log2BlockSize = 9; // block_size = 512 + + int blockSize, numBlocks, bufferSize; + _computeSizes(n, log2BlockSize, &blockSize, &numBlocks, &bufferSize); + + int *buffer; + cudaMalloc(&buffer, sizeof(int) * bufferSize); + + cudaMemcpy(buffer, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + // if integer overflow on the GPU were well-defined we would be able to get away without zeroing the rest + cudaMemset(buffer + n, 0, sizeof(int) * (bufferSize - n)); + timer().startGpuTimer(); - // TODO + kernScanPerBlock<<>>(buffer, nullptr); timer().endGpuTimer(); + + odata[0] = 0; + cudaMemcpy(odata, buffer, sizeof(int) * n, cudaMemcpyDeviceToHost); + + cudaFree(buffer); + } + + __global__ void kernConvertToBinary(int n, int *odata, const int *idata) { + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + if (iSelf >= n) { + return; + } + odata[iSelf] = idata[iSelf] != 0 ? 1 : 0; + } + + __global__ void kernCompact(int n, int *out, const int *data, const int *accum) { + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + if (iSelf >= n) { + return; + } + + int val = data[iSelf]; + if (val != 0) { + out[accum[iSelf]] = val; + } } /** @@ -31,10 +122,32 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { + int log2BlockSize = 9; // block_size = 512 + + int blockSize, numBlocks, bufferSize; + _computeSizes(n, log2BlockSize, &blockSize, &numBlocks, &bufferSize); + + int *data, *accum, *out; + cudaMalloc(&data, sizeof(int) * bufferSize); + cudaMalloc(&accum, sizeof(int) * bufferSize); + cudaMalloc(&out, sizeof(int) * bufferSize); + + cudaMemcpy(data, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + cudaMemset(data + n, 0, sizeof(int) * (bufferSize - n)); + timer().startGpuTimer(); - // TODO + kernConvertToBinary<<>>(n, accum, data); + kernScanPerBlock<<>>(accum, nullptr); + kernCompact<<>>(n, out, data, accum); timer().endGpuTimer(); - return -1; + + int res; + cudaMemcpy(odata, out, sizeof(int) * n, cudaMemcpyDeviceToHost); + cudaMemcpy(&res, accum + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + if (idata[n - 1] != 0) { + ++res; + } + return res; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..e003bf5 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -11,15 +11,48 @@ namespace StreamCompaction { static PerformanceTimer timer; return timer; } - // TODO: __global__ + + constexpr int block_size = 128; + + __global__ void kernScanPass(int n, int diff, int *odata, const int *idata) { + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + if (iSelf > n) { + return; + } + if (iSelf >= diff) { + odata[iSelf] = idata[iSelf] + idata[iSelf - diff]; + } else { + odata[iSelf] = idata[iSelf]; + } + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); + int *buf1 = nullptr, *buf2 = nullptr; + cudaMalloc(&buf1, sizeof(int) * n); + cudaMalloc(&buf2, sizeof(int) * n); + + cudaMemcpy(buf1, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + + { + timer().startGpuTimer(); + + int num_blocks = (n + block_size - 1) / block_size; + for (int diff = 1; diff < n; diff *= 2) { + kernScanPass<<>>(n, diff, buf2, buf1); + std::swap(buf1, buf2); + } + + timer().endGpuTimer(); + } + + odata[0] = 0; + cudaMemcpy(odata + 1, buf1, sizeof(int) * (n - 1), cudaMemcpyDeviceToHost); + + cudaFree(buf1); + cudaFree(buf2); } } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..0ee6ce5 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -18,11 +18,16 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + thrust::device_vector dev_vec(idata, idata + n), dev_out(n); + timer().startGpuTimer(); - // TODO use `thrust::exclusive_scan` + // use `thrust::exclusive_scan` // example: for device_vectors dv_in and dv_out: // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + thrust::exclusive_scan(dev_vec.begin(), dev_vec.end(), dev_out.begin()); timer().endGpuTimer(); + + thrust::copy(dev_out.begin(), dev_out.end(), odata); } } } From a4d3c9e20b32e4b69213def39b3f533addf9cdc9 Mon Sep 17 00:00:00 2001 From: Xuanyi Zhou Date: Fri, 18 Sep 2020 23:42:59 -0400 Subject: [PATCH 2/4] support arbitrary length array for efficient implementation --- src/main.cpp | 2 +- stream_compaction/common.cu | 16 ++++++- stream_compaction/efficient.cu | 85 ++++++++++++++++++++-------------- stream_compaction/efficient.h | 3 ++ 4 files changed, 68 insertions(+), 38 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 896ac2b..e7ecfd7 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1 << 27; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..3ede5e4 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -23,7 +23,11 @@ namespace StreamCompaction { * which map to 0 will be removed, and elements which map to 1 will be kept. */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { - // TODO + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + if (iSelf >= n) { + return; + } + bools[iSelf] = idata[iSelf] != 0 ? 1 : 0; } /** @@ -32,7 +36,15 @@ namespace StreamCompaction { */ __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { - // TODO + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + if (iSelf >= n) { + return; + } + + int val = bools[iSelf]; + if (val != 0) { + odata[indices[iSelf]] = idata[iSelf]; + } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2ab90b5..48d1749 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -1,7 +1,11 @@ +#include "efficient.h" + +#include + #include #include + #include "common.h" -#include "efficient.h" namespace StreamCompaction { namespace Efficient { @@ -11,6 +15,14 @@ namespace StreamCompaction { return timer; } + constexpr int log2BlockSize = 9; // 512 + constexpr int blockSize = 1 << log2BlockSize; + + constexpr int numBanks = 32; + /*__host__ __device__ int _transformIndex(int i) { + + }*/ + __global__ void kernScanPerBlock(int *data, int *lastData) { extern __shared__ int buffer[]; data += blockIdx.x * blockDim.x; @@ -57,23 +69,44 @@ namespace StreamCompaction { } } - void _computeSizes(int n, int log2BlockSize, int *blockSize, int *numBlocks, int *bufferSize) { - *blockSize = 1 << log2BlockSize; + __global__ void kernAddConstantToBlock(int *data, const int *amount) { + data[blockIdx.x * blockDim.x + threadIdx.x] += amount[blockIdx.x]; + } + + void _computeSizes(int n, int log2BlockSize, int *numBlocks, int *bufferSize) { *numBlocks = n >> log2BlockSize; - if ((n & (*blockSize - 1)) != 0) { + if ((n & ((1 << log2BlockSize) - 1)) != 0) { ++*numBlocks; } *bufferSize = *numBlocks << log2BlockSize; } + void dev_scan(int n, int *dev_data) { + assert((n & (blockSize - 1)) == 0); + + if (n > blockSize) { + int numBlocks = n >> log2BlockSize, numIndirectBlocks, indirectSize; + _computeSizes(numBlocks, log2BlockSize, &numIndirectBlocks, &indirectSize); + + int *buffer; + cudaMalloc(&buffer, sizeof(int) * indirectSize); + + kernScanPerBlock<<>>(dev_data, buffer); + dev_scan(indirectSize, buffer); + kernAddConstantToBlock<<>>(dev_data, buffer); + + cudaFree(buffer); + } else { // just scan the block + kernScanPerBlock<<<1, blockSize, blockSize * sizeof(int)>>>(dev_data, nullptr); + } + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - int log2BlockSize = 9; // block_size = 512 - - int blockSize, numBlocks, bufferSize; - _computeSizes(n, log2BlockSize, &blockSize, &numBlocks, &bufferSize); + int numBlocks, bufferSize; + _computeSizes(n, log2BlockSize, &numBlocks, &bufferSize); int *buffer; cudaMalloc(&buffer, sizeof(int) * bufferSize); @@ -83,33 +116,14 @@ namespace StreamCompaction { cudaMemset(buffer + n, 0, sizeof(int) * (bufferSize - n)); timer().startGpuTimer(); - kernScanPerBlock<<>>(buffer, nullptr); + dev_scan(bufferSize, buffer); timer().endGpuTimer(); odata[0] = 0; cudaMemcpy(odata, buffer, sizeof(int) * n, cudaMemcpyDeviceToHost); cudaFree(buffer); - } - - __global__ void kernConvertToBinary(int n, int *odata, const int *idata) { - int iSelf = blockIdx.x * blockDim.x + threadIdx.x; - if (iSelf >= n) { - return; - } - odata[iSelf] = idata[iSelf] != 0 ? 1 : 0; - } - - __global__ void kernCompact(int n, int *out, const int *data, const int *accum) { - int iSelf = blockIdx.x * blockDim.x + threadIdx.x; - if (iSelf >= n) { - return; - } - - int val = data[iSelf]; - if (val != 0) { - out[accum[iSelf]] = val; - } + checkCUDAError("efficient scan"); } /** @@ -122,10 +136,10 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - int log2BlockSize = 9; // block_size = 512 + constexpr int log2BlockSize = 9; // block_size = 512 - int blockSize, numBlocks, bufferSize; - _computeSizes(n, log2BlockSize, &blockSize, &numBlocks, &bufferSize); + int numBlocks, bufferSize; + _computeSizes(n, log2BlockSize, &numBlocks, &bufferSize); int *data, *accum, *out; cudaMalloc(&data, sizeof(int) * bufferSize); @@ -136,14 +150,15 @@ namespace StreamCompaction { cudaMemset(data + n, 0, sizeof(int) * (bufferSize - n)); timer().startGpuTimer(); - kernConvertToBinary<<>>(n, accum, data); - kernScanPerBlock<<>>(accum, nullptr); - kernCompact<<>>(n, out, data, accum); + Common::kernMapToBoolean<<>>(n, accum, data); + dev_scan(bufferSize, accum); + Common::kernScatter<<>>(n, out, data, data, accum); timer().endGpuTimer(); int res; cudaMemcpy(odata, out, sizeof(int) * n, cudaMemcpyDeviceToHost); cudaMemcpy(&res, accum + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + checkCUDAError("efficient compaction"); if (idata[n - 1] != 0) { ++res; } diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h index 803cb4f..f11a236 100644 --- a/stream_compaction/efficient.h +++ b/stream_compaction/efficient.h @@ -6,6 +6,9 @@ namespace StreamCompaction { namespace Efficient { StreamCompaction::Common::PerformanceTimer& timer(); + /// Scans a device array. The input size must be a multiple of 512. + void dev_scan(int n, int *dev_data); + void scan(int n, int *odata, const int *idata); int compact(int n, int *odata, const int *idata); From 3603df2076434a2674f5a076fbcf389cc87f90d9 Mon Sep 17 00:00:00 2001 From: Xuanyi Zhou Date: Sun, 20 Sep 2020 01:47:30 -0400 Subject: [PATCH 3/4] improved scan implementation, radix sort --- src/main.cpp | 33 +++++++++- stream_compaction/efficient.cu | 117 ++++++++++++++++++++++++++------- stream_compaction/efficient.h | 2 + 3 files changed, 127 insertions(+), 25 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index e7ecfd7..6779123 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -7,13 +7,14 @@ */ #include +#include #include #include #include #include #include "testing_helpers.hpp" -const int SIZE = 1 << 27; // feel free to change the size of array +const int SIZE = 1 << 25; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; @@ -147,6 +148,36 @@ int main(int argc, char* argv[]) { //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); + + printf("\n"); + printf("*********************\n"); + printf("** RADIX SORT TEST **\n"); + printf("*********************\n"); + + // here we write our own version of genArray that does not make use of grandpa's functions + // because on my machine RAND_MAX is 0x7FFF which means not all bits can be tested + std::default_random_engine rand(std::random_device{}()); + std::uniform_int_distribution dist(0, 2000000000); + for (std::size_t i = 0; i < SIZE; ++i) { + a[i] = dist(rand); + } + printArray(SIZE, a, true); + + printDesc("radix sort, power-of-two"); + std::memcpy(b, a, sizeof(int) * SIZE); + std::sort(b, b + SIZE); + StreamCompaction::Efficient::radix_sort(SIZE, c, a); + printCmpResult(SIZE, b, c); + printArray(SIZE, c, true); + + printDesc("radix sort, non-power-of-two"); + std::memcpy(b, a, sizeof(int) * NPOT); + std::sort(b, b + NPOT); + StreamCompaction::Efficient::radix_sort(NPOT, c, a); + printCmpResult(NPOT, b, c); + printArray(NPOT, c, true); + + system("pause"); // stop Win32 console from closing on exit delete[] a; delete[] b; diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 48d1749..c69e76a 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -15,57 +15,56 @@ namespace StreamCompaction { return timer; } + // block size when processing. cuda block size is half this constexpr int log2BlockSize = 9; // 512 constexpr int blockSize = 1 << log2BlockSize; - constexpr int numBanks = 32; - /*__host__ __device__ int _transformIndex(int i) { - - }*/ - __global__ void kernScanPerBlock(int *data, int *lastData) { extern __shared__ int buffer[]; - data += blockIdx.x * blockDim.x; + data += blockIdx.x * blockDim.x * 2; // copy data to shared memory - buffer[threadIdx.x] = data[threadIdx.x]; + buffer[threadIdx.x * 2] = data[threadIdx.x * 2]; + buffer[threadIdx.x * 2 + 1] = data[threadIdx.x * 2 + 1]; __syncthreads(); int lastElem = 0; if (lastData && threadIdx.x == blockDim.x - 1) { - lastElem = buffer[threadIdx.x]; + lastElem = buffer[threadIdx.x * 2 + 1]; } // upward pass - for (int gap = 2; gap < blockDim.x; gap <<= 1) { - if ((threadIdx.x & (gap - 1)) == gap - 1) { - buffer[threadIdx.x] += buffer[threadIdx.x - (gap >> 1)]; + for (int halfGap = 1; halfGap < blockDim.x; halfGap <<= 1) { + if (threadIdx.x < blockDim.x / halfGap) { + buffer[(threadIdx.x * 2 + 2) * halfGap - 1] += buffer[(threadIdx.x * 2 + 1) * halfGap - 1]; } __syncthreads(); } if (threadIdx.x == blockDim.x - 1) { - int halfIdx = threadIdx.x >> 1; - buffer[threadIdx.x] = buffer[halfIdx]; - buffer[halfIdx] = 0; + int halfIdx = blockDim.x - 1; + buffer[blockDim.x * 2 - 1] = buffer[threadIdx.x]; + buffer[threadIdx.x] = 0; } __syncthreads(); // downward pass - for (int gap = blockDim.x >> 1; gap > 1; gap >>= 1) { - if ((threadIdx.x & (gap - 1)) == gap - 1) { - int prevIdx = threadIdx.x - (gap >> 1); - int sum = buffer[threadIdx.x] + buffer[prevIdx]; - buffer[prevIdx] = buffer[threadIdx.x]; - buffer[threadIdx.x] = sum; + for (int halfGap = blockDim.x >> 1; halfGap >= 1; halfGap >>= 1) { + if (threadIdx.x < blockDim.x / halfGap) { + int prevIdx = (threadIdx.x * 2 + 1) * halfGap - 1; + int thisIdx = prevIdx + halfGap; + int sum = buffer[thisIdx] + buffer[prevIdx]; + buffer[prevIdx] = buffer[thisIdx]; + buffer[thisIdx] = sum; } __syncthreads(); } // copy data back - data[threadIdx.x] = buffer[threadIdx.x]; + data[threadIdx.x * 2] = buffer[threadIdx.x * 2]; + data[threadIdx.x * 2 + 1] = buffer[threadIdx.x * 2 + 1]; if (lastData && threadIdx.x == blockDim.x - 1) { - lastData[blockIdx.x] = lastElem + buffer[threadIdx.x]; + lastData[blockIdx.x] = lastElem + buffer[threadIdx.x * 2 + 1]; } } @@ -91,16 +90,17 @@ namespace StreamCompaction { int *buffer; cudaMalloc(&buffer, sizeof(int) * indirectSize); - kernScanPerBlock<<>>(dev_data, buffer); + kernScanPerBlock<<>>(dev_data, buffer); dev_scan(indirectSize, buffer); kernAddConstantToBlock<<>>(dev_data, buffer); cudaFree(buffer); } else { // just scan the block - kernScanPerBlock<<<1, blockSize, blockSize * sizeof(int)>>>(dev_data, nullptr); + kernScanPerBlock<<<1, blockSize / 2, blockSize * sizeof(int)>>>(dev_data, nullptr); } } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ @@ -126,6 +126,7 @@ namespace StreamCompaction { checkCUDAError("efficient scan"); } + /** * Performs stream compaction on idata, storing the result into odata. * All zeroes are discarded. @@ -164,5 +165,73 @@ namespace StreamCompaction { } return res; } + + + __global__ void kernExtractBit(int n, int bit, int *odata, const int *idata) { + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + if (iSelf >= n) { + return; + } + odata[iSelf] = (idata[iSelf] & (1 << bit)) != 0 ? 1 : 0; + } + + __global__ void kernNegate(int *odata, const int *idata) { + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + odata[iSelf] = idata[iSelf] == 0 ? 1 : 0; + } + + __global__ void kernRadixSortScatter( + int n, int numFalses, int bit, int *odata, const int *idata, const int *trues, const int *falses + ) { + int iSelf = blockIdx.x * blockDim.x + threadIdx.x; + if (iSelf >= n) { + return; + } + int value = idata[iSelf], index; + if ((value & (1 << bit)) != 0) { + index = trues[iSelf] + numFalses; + } else { + index = falses[iSelf]; + } + odata[index] = value; + } + + void radix_sort(int n, int *odata, const int *idata) { + constexpr int numIntBits = sizeof(int) * 8 - 1; + + int numBlocks, bufferSize; + _computeSizes(n, log2BlockSize, &numBlocks, &bufferSize); + + int *data1, *data2, *trues, *falses; + cudaMalloc(&data1, sizeof(int) * n); + cudaMalloc(&data2, sizeof(int) * n); + cudaMalloc(&trues, sizeof(int) * bufferSize); + cudaMalloc(&falses, sizeof(int) * bufferSize); + + cudaMemcpy(data1, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + + for (int i = 0; i < numIntBits; ++i) { + kernExtractBit<<>>(n, i, trues, data1); + kernNegate<<>>(falses, trues); + dev_scan(bufferSize, trues); + dev_scan(bufferSize, falses); + int numFalses, lastElem; + cudaMemcpy(&lastElem, data1 + (n - 1), sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(&numFalses, falses + (n - 1), sizeof(int), cudaMemcpyDeviceToHost); + if ((lastElem & (1 << i)) == 0) { + ++numFalses; + } + kernRadixSortScatter<<>>(n, numFalses, i, data2, data1, trues, falses); + std::swap(data1, data2); + } + + cudaMemcpy(odata, data1, sizeof(int) * n, cudaMemcpyDeviceToHost); + + cudaFree(data1); + cudaFree(data2); + cudaFree(trues); + cudaFree(falses); + checkCUDAError("radix sort"); + } } } diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h index f11a236..f7706d8 100644 --- a/stream_compaction/efficient.h +++ b/stream_compaction/efficient.h @@ -12,5 +12,7 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata); int compact(int n, int *odata, const int *idata); + + void radix_sort(int n, int *odata, const int *idata); } } From 616d4fd0487089291a33ca0c340ca99248b2bea5 Mon Sep 17 00:00:00 2001 From: Xuanyi Zhou Date: Tue, 22 Sep 2020 21:47:36 -0400 Subject: [PATCH 4/4] optimized efficient implementation, thrust compact implementation, readme --- README.md | 122 +++++++++++++++++++++++++++++++-- data.xlsx | Bin 0 -> 15552 bytes img/perf.png | Bin 0 -> 74128 bytes src/main.cpp | 20 +++++- stream_compaction/common.cu | 3 +- stream_compaction/efficient.cu | 72 ++++++++++++------- stream_compaction/naive.cu | 3 +- stream_compaction/thrust.cu | 19 +++++ stream_compaction/thrust.h | 2 + 9 files changed, 207 insertions(+), 34 deletions(-) create mode 100644 data.xlsx create mode 100644 img/perf.png diff --git a/README.md b/README.md index 0e38ddb..378353b 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,122 @@ CUDA Stream Compaction **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2** -* (TODO) YOUR NAME HERE - * (TODO) [LinkedIn](), [personal website](), [twitter](), etc. -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Xuanyi Zhou + * [LinkedIn](https://www.linkedin.com/in/xuanyi-zhou-661365192/), [Github](https://github.com/lukedan) +* Tested on: Windows 10, i7-9750H @ 2.60GHz 32GB, RTX 2060 6GB -### (TODO: Your README) +## Features -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +- Regular CPU and thrust implementations of scan and compact operations. +- Naive implementation of the scan operation. +- Recursive efficient scan implementation that supports arbitrary array lengths, and makes use of shared memory & bank conflict optimizations. +- Efficient compact implementaiton making use of the efficient scan operation. +## Questions + +- Roughly optimize the block sizes of each of your implementations for minimal run time on your GPU. + + For Naive scan, changing the block size does not affect executing time much as the operation is bounded by memory bandwidth. For efficient scan and compaction, the optimal block size is around 64/128. From tracing the program it can be seen that thrust also uses a block size of 128. + +- Compare all of these GPU Scan implementations (Naive, Work-Efficient, and Thrust) to the serial CPU version of Scan. Plot a graph of the comparison (with array size on the independent axis). + + Below is the graph of run time of all implementations: + + ![](img/perf.png) + + See `data.xlsx` for the raw data. Note that since the efficient scan implementation is recursive, some memory allocations are included in the final time. + + For small array sizes, all methods are fairly fast and the run times are relatively unpredictable. As the array size increases, the relationship between the run times become consistent: the naive scan is the slowest, followed by the CPU implementation, the work-efficient scan, and finally the thrust implementation is the fastest. + + Looking at the Nsight timeline and kernel calls it can be seen that the number of threads spawned by the thrust implementation is far less than the work-efficient implementation which requires one thread for every two array elements. The thrust implementation also uses far more registers per thread. This suggests that the thrust implementation may be using an optimization mentioned in the GPU gems chapter, which is performing a scan operation for a few elements in each thread in serial, then aggregating the results by performing a scan on the sums. + + Interestingly, the time gaps between copying data from CPU to GPU and back are roughly the same for the work-efficient implementation and the thrust implementation as shown in the timeline, but the work-efficient implementation is busy throughout the period while the thrust implementation is largely idle for the first half of the gap. Currently I do not have a reasonable explanation for this. + +- Write a brief explanation of the phenomena you see here. + + Since the naive implementation does not use any shared memory, it's understandable that it would be slower than the CPU implementation as the GPU has a much lower clock speed than the CPU, and does not have nearly the same amount of cache. + + The thrust implementation and the work-efficient implementation are faster than the CPU version. The thrust version is about three times faster as it may have utilized other optimizations such as the one mentioned in the previous answer. + + The reason why the optimization mentioned before works may be that it reduces the number of processed elements by a large factor (4x if 4 elements are summed in serial) while summing a few numbers in serial is relatively cheap and is very well parallelized. + +- Paste the output of the test program into a triple-backtick block in your README. + + The tests for radix sort can be seen at the end of the output. + + ``` + **************** + ** SCAN TESTS ** + **************** + [ 45 4 11 40 10 5 35 48 33 44 0 28 24 ... 41 0 ] + ==== cpu scan, power-of-two ==== + elapsed time: 78.9076ms (std::chrono Measured) + [ 0 45 49 60 100 110 115 150 198 231 275 275 303 ... -1007647599 -1007647558 ] + ==== cpu scan, non-power-of-two ==== + elapsed time: 77.9638ms (std::chrono Measured) + [ 0 45 49 60 100 110 115 150 198 231 275 275 303 ... -1007647704 -1007647672 ] + passed + ==== naive scan, power-of-two ==== + elapsed time: 133.709ms (CUDA Measured) + passed + ==== naive scan, non-power-of-two ==== + elapsed time: 112.619ms (CUDA Measured) + passed + ==== work-efficient scan, power-of-two ==== + elapsed time: 11.0549ms (CUDA Measured) + passed + ==== work-efficient scan, non-power-of-two ==== + elapsed time: 11.0906ms (CUDA Measured) + passed + ==== thrust scan, power-of-two ==== + elapsed time: 4.04538ms (CUDA Measured) + passed + ==== thrust scan, non-power-of-two ==== + elapsed time: 4.0943ms (CUDA Measured) + passed + + ***************************** + ** STREAM COMPACTION TESTS ** + ***************************** + [ 3 0 2 2 1 0 0 2 2 3 2 3 1 ... 2 0 ] + ==== cpu compact without scan, power-of-two ==== + elapsed time: 307.968ms (std::chrono Measured) + [ 3 2 2 1 2 2 3 2 3 1 3 2 2 ... 2 2 ] + passed + ==== cpu compact without scan, non-power-of-two ==== + elapsed time: 304.929ms (std::chrono Measured) + [ 3 2 2 1 2 2 3 2 3 1 3 2 2 ... 2 2 ] + passed + ==== cpu compact with scan ==== + elapsed time: 422.665ms (std::chrono Measured) + [ 3 2 2 1 2 2 3 2 3 1 3 2 2 ... 2 2 ] + passed + ==== work-efficient compact, power-of-two ==== + elapsed time: 25.4493ms (CUDA Measured) + [ 3 2 2 1 2 2 3 2 3 1 3 2 2 ... 2 2 ] + passed + ==== work-efficient compact, non-power-of-two ==== + elapsed time: 23.5873ms (CUDA Measured) + [ 3 2 2 1 2 2 3 2 3 1 3 2 2 ... 2 2 ] + passed + ==== thrust compact, power-of-two ==== + elapsed time: 4.44602ms (CUDA Measured) + passed + ==== thrust compact, non-power-of-two ==== + elapsed time: 4.03315ms (CUDA Measured) + passed + + ********************* + ** RADIX SORT TEST ** + ********************* + [ 943102324 1728649027 1523795418 230368144 1853983028 219035492 1373487995 539655339 345004302 1682352720 528619710 1142157171 735013686 ... 646714987 484939495 ] + ==== radix sort, power-of-two ==== + elapsed time: 1305.98ms (CUDA Measured) + passed + [ 42 45 55 58 63 67 78 122 162 170 188 206 221 ... 1999999985 1999999998 ] + ==== radix sort, non-power-of-two ==== + elapsed time: 1338.5ms (CUDA Measured) + passed + [ 42 45 55 58 63 67 78 122 162 170 188 206 221 ... 1999999985 1999999998 ] + Press any key to continue . . . + ``` diff --git a/data.xlsx b/data.xlsx new file mode 100644 index 0000000000000000000000000000000000000000..055bd1b30c3dfd2e2211a0f82b679582ce0eea0d GIT binary patch literal 15552 zcmeHu^Lu35wsve=9Xsi;W83c7wr$(CZ9D1Mw$Wk79jlXk>Am+ky-%NgpYIR2cl}VS zs%DM%sWI0Wb6~zxRst9V1po{H0ssJj0KkrmiZcli06-QD000>P0!UrJ+RDMu%0c_H ztBs+(CXI`w1%4g~5P1#&(EIQI@A_Xn0~Lv5(tUIYLieIS1ol3Puh;QHFrE1gAyLY4 zcQ412mg%eG-rwF!t~cRJeqMre!hYKpx$~N}I(jqN)rKUl#P*h8UkeryT5{>yN zZ_uQ_tO`(onz8M;nOth;1@ey;3|E{WE^N~~^gs>|kf+V=-CH|oxgy{Y-EzQZX?K(} z!sK3mH~1{dd}rXTMh*{h+QmrF%}tH0d!-V8su6u&+~Ca=hlBiW*VBi~Dcz-%^J9x* z+}+9btcAs09t;LBX7D4FkYlgo%KT4i@ZLNF&uB}6#Zdv>!BB8g=PLEA2&L{?8qCu- z#-PU#=|gWl56}PrZ*RZ=vj5@Ejf!+cH}5to`R=XI@9wN^XJ}zhOY`&f|J?Y0vG4v( z^@=ztnLaw0pbOC-!Nd0p8!-s{63%?0?F65_e8kq_8zb|Hu{XNNa1cIY`2&l1cY6I8 zUEAP{JRT;v+xb)xj)cNN)Z|B7t2u?|Amn3Xk@wFe(Y2kL^E>%R_mE5T_nzFp5 zBv)c&i%59xQm__bibfd=61fm32#qV%S8Y&AbyM%53UGm6;jA*KrkN$@IDRVKdnu{t z2$m;=P4aX$6?NEN&v>QEW7vY=_7O`_!Gz7UMlZ{bozPX=z@qzFFue=;#p9!N#;`o$ zC!`04DUl(PoLg_ThEJDcnXdgTkbULD-$(qy@k%z|BinyFiF2tf1~4!Hz|DKigMXhH z7YkZvYddp2YiskL@vcHqQ)Z10#WTC%P3Mt|SUSET62hdxupC4|=+UJ#AYzvpS1e-% z?j!$?drh(c`DI{#p~|i*dyn(ZL-yz3bV3nl4O4pbjp;fs!)nc zCs|p9&Y!Z-TR~AttPu5Xaba3DpazHH7C>pGpz~a$VW*c;OB{lgYVNyHKW{81_8G_l z_D2o}wp#PjQ>t)kCd$vG0IkwC_A((jr7n>l+)nUW`fyB^6q)Bs;W%O`I3GiGnj9h&;z+ny!_j64G9A`!c!$9vXxgL(Yr#$j-9Na)h%jt zs#FuVdY>|bz7E4%!Ix)^E4Z()(5XDZ9)DvP7bnr^Z>G+I>#Z32|?sp_KjjM;HFBVNYCzt4wB5v zd+$D@XU29LB)?}F6Yi=1(v4{L1bGYlE?>k(ScpEpfq5k!i`v1@|2u7wLfa8_f4F(9 zl_{r-Otk*9vXrrV1KgVXh5Yx%&ErAonFH~JQVp_YKr4)m3+biEOYs6y;IFWq08I_% zJS{4Q-6;XU?GU%RsU$AJGcAS=Fwy$yAB4}L#V25)=hJW#r4+}{DkZiln>~>YO8l0y za*HTT?kX`?@blJEtMR1hz0ixVA7s z^}Uhq$aY8MOs@625H8u!DH;MSn1Fc0bwv z=?(wwFTUIf$=B~AZSO$^3jhM}-CzEUD}VQzf5jNU_f+xS>;89N74lLN{d5SO(67OC z&go95h${}Xgr|xp@DL*n6lVn8Eo)gn~D% zvq7KCCg~g*%jJ~h2V-0&gPXP+2m3|G)&G=~)M*DBA0&90**?z8gGuN1qVw|w9 z05jj_38)@=gzg}&I(`^QxlWq*8zxZp3wKp*|bg-xW<3|5; zt<6wTw=QQz@X|JY1K#NzWT6uji<0Iqt8AdEynB)ijn4^+0ZnUAJM-S6KhRr`56dV4 zWw>@7%r@qXV0^nu!pjV&F0WwysNq+0GOVCxQ9tMP_4RRRSKu4FVrE9uNG~K6f~$PrlHG;kp*-oDk*wDZ z`Jity*~-WQn$bQA_JCD#Ap_0ql17%9WkJGFLC~Xe@TrF5;*~>KD<_4v))c&HY<5q- zz0>&d{b0$--qES4H|6NwBSgm*gpB-QWdIMkWFX+wL;4R3GJ$;39 zFNZXTYcPZrDr|-8AJ?l5nh3s}d_0r_dPy8HH{|4Sy!z43Iogo49nd1JncuPxlnzjpbXMsGxL9%r)4h39B|J3x z;+d+_t$q=8tMrrFg3xR8A|g|@7j-~6CiJ}%wZ(d!+#9rB4c?;utvVh`aqAr?2o5wO+x%z&{z7A$MMFZp#;?t%ktMxR>Cy{JpDXbz8!(6_3FCPUK90C`>k?Lc zL~s_T44`9NJ7C$iF*(ggq#&1=69a{%FjGeUI7kjPK<;%HbZSEo3Q{W}QEcF=a}y79 zW|KfaXg!Z<3r~?NrW)QntLK-4%geareVTEq189);fvmD`_>h2kEK!q<1IL(9C#YUG zRG2_P>SXUDx|!3WS#?7SL1IvL&55qA;7fWqf;q;`QpCd-MaD)rhAa>QMmw+J2Shd( zHW@N>Q@cgCs)l8e6uWvowLr&8T_7C-MjJHhMLItZM0ujlvxy>I%H;@Q0n&(JL{^{m zOCV*bbwjaAwJ$(a;?fX#f-0kq#e3uky=9cP&)$@i5Ps9_y=$uLI+>00!$c` z+$utn;vpbOBi7%h6HRdX$U{7?_syNjKf(oK0NJ+RIeQ@vtW&G`Z`tmhn`w%!6)bm3 zP#xmfN;Fz20tP$uKisuf_5t#sVXTeRH6IoEn246vYOyqZXf0O&bcj!NDxVz{2x2w< zDur}{CSnrc;?TZ0^Z|}iLX^ifwwd+Y1Sa%yvmZCyfur>)sh%-dY2^-Zr6v>_ss z>Drlel{L0=8b1(Ay} zt}^wOHGQn8fy7*;V8(gU!+vsx1zoelfwpK8QJPIxY^_qOqy z4^*5qqzyUov)5U6EixeGB4Nw}W+Fse$3COGoJtc(8R5@liz4pCf-w4a`hyC*w|{wYkMgtPO0hs;-IL=EN$Q#X7~{XTwOvS7u~ zZ3*1jc{curGkzXqC(i!+fD_q%P)?I})Z4=(XWFD0X)D12H1!FqoL9T{bsaku1<#`T znKp%~>&~3|yJ~>X7CTRoUeoZ4uP6YKLv*td;pcM~O>JB}LuNUrvpele-xkWjw63+c zw%0HBr@EKf&WXn`2Gw$AEsdD?qy)_%R*)epay4ijeg5Y-#cG=Erb9nr(c3FN^ov^JuTK%l3V z1^M!^>Y0$iXDKOU)kf1I=(`gXr9QBh);d!QuHX?rg2a?^{xZNhqi{QQBiMG?sYLj9 zo7nNw%wp&d&d9k$>RSZj5g_;I_A$zO_L#!UIj;VBf;pJdH8y*UP9e314$cttJ8d$9 z6uTCb865pbgnF4d{^ zKOxAv?*%-uJs52S7KV20VEmS^g$JVfxeNW=^5W`o%rSQwe{(=hvvPg668fGZ_lB36 zVsqE*Y*JS%4JJpNXYQkJ$CmL!sG&HFruf#G%w541Cvy+PM-Vxt@rM276QyhS% z&mlXZ;|fm1HKSFMldVQ>>U5=AfHh$%ioCH$NqcRME_M#8rJce>ut;jHsl*Ko5z`=t z3#E7&8OBxWLgKJ)3@GKzxj8c?U!pZ65ZMZ0Kg}fikxtsK!5(vGL`KeZoJs3$!6#qh zm<5l4(&jO;a$<*na{sy|ta?C!&efk!y)Ev*z!VbNT`C~9?zI1XX@M8HkRVLG5zOy< zHOUf?>SJ@Pe?hv~;2Z@mG}oD0c!Yx!G*q@};B6^`w|h;40kr5n-zPRFVTagFL0LVo zd`e(e0-PxCGSyFHAd`$s{7Um%rJ)qbnlHR^3$lw*-#GanBl5ys!ITNxit^7F>8H=H z=mmIL1X$BXAn!7bH|mScFxjL&g|p4cqOzH<4&1y*U162$Au*{K&Q>lVG>C$z?=)o*Rm8firQ)4RD_3UVE1ifd zJ{m)xjjHHen!Xp;-!XV{x#}47*m3wyhD?^wF&iH6fGz6^y^Ly)mjA5?1 zJ{AShDnGQ2_nWFwOvehL~xdwgLVmMCVBg@>UewYSI(6r7Vc9v+B;z!%a9LK(# z23s4c+fer!gF+0~M-iJq#feXC@W_K7mv-0$DgF`r_#qAY{EiDvxK3l;I)T~YqmLQ7 zlu2ltiJ9Jfs(#$&Ghy^z;f&JKIh0LkmAO`3vcsW)w27-Q(B2Y(5G3So_$tS>cmX86 z#ZadQtBiSNhESnZA~3ygiFMD8>d{5n_v{Y)kWqza;Ik3{nEYwG37 zUCbdMZn} zQBRygHl`ROX_Z{ANDsgJ{ch}?n?648Ka!7De`Gt z`|P2)@A^oAe6}Bidkn;ldo0{CL4=S7%8v|OlTk)O+^s}D z6~GeDw7X%@`C@X;@DzqE7+R@?G9WI32k|-84GCI8mMc&_I-SQU$e(cs%{YGm^QB0; zeg$vH?66DQ8D~vlbctfoksLTK;v{#3>#g36bNv3AIofd(q+`J4ff{l@WWu)0V>rX^8v zJHn0$$7T_{+3?;3*nM2Vg(s5P15t2&LGjYVGYjH9!r>vg8WjTWz~<~e^yh0$8AB#n z8v2X-z494jE&(d`S1S**yt7L&0xk1PZfY7Qq~CUX!yyc_QH8TVOV2qRxid*=Qpw!H zA>p-W%_A{FVUWtBOkelO@N9zMaSH97-{A9Yq0ssIV1EX5{*Z*#4;EzaHp7#6hvt&t z37CDtzCv@y1ZAcVO1_$>(=Jk;BYm*Kt-9E5ko85yA>*Jwq;eu-gE@ZK_9=P1g2j_^ z<`@%^8pU!AnLRQRGYun9+GuIt{#oCun=mX1o5 zz|zZ&Q!~bGGt1radgpF`?`ed`_(EZ#QHG5bJ4iB6TBmEB46t1x0A{AaRoiFeVJiKu zTcbcGT#AyYdq3tf`7ic&H!@&A=qq z>C4cU2ZPcv$_zL{*1?O=j5nmsBqs8E8%`V!7jy~ zVw@J7g1#mo-(WaUmRTfJjQQMKOCCoEKlRnw zxr})w)fu4k35g_RxMU0Z4EifT_Nn?@{k2zFCq3p5F;n00$xL(C>OEJJuU*ro#Zsy5 z)ohrcgc*DJfNmYt7TTIY9>-m$^|MGTD#qBF;~4PDCT|LI824b6>)F8wCLoH*>oj+h z3`B}7;e>RB!vZOD2DUc}DGm1PQf5>&>VT5+6?tTYcOw*(c$IBcZ7>{(sy%2!4+A!+ z%K>4U{lui_sQ2&fRu3*EzHxk^;lmwNfAUIYP{T`wv|A#y!J`7;zz7m3R#i@+wi99M znkq>{sQW;__Vu{c4#JpttcbW3vCdRcOHp$R9k#s}p|Sug6pwQsGm+=KpG6}6WH;I1 z(?{cb^l=>$^l+VdMZ2WBwCwmaZ3Hm4U`bujAFv)1^;dndV)ZZ%1fZ4(`}l!zo}VJ& zBR-25y?BXXFu}?b%fcY-N^1AxRK=zcVHz@Zp(rGyCz6-aq&zU@UpLdGJrR9&w8WVD z97UP9HAYkqpnz;NG>c8wn%9xgCAkZh0FdVMb(muKSV4Y{8nwu@&n-?))9C1wdV;}~ z;523sQ(^YdS#)tE>9O3dXw+jW=zQKK`7!-b165y>Gl0gCw#FtC9HO2A#!GvK?G#w# zpzBI6#&+QYNhWqe3AWfqYwWx{MT5?sK9-!AD?LI|4WeHeOByQBDuEUhPfA!0)5f)#UcpZt z!GxtuXki8S40U_^i7hG95q7I8R=}i+N>}vxf`DV5ttgbh7{KoQ)Cg)C)BqZkyqblP zP*W?CUnf86-rJ@oVy6IOiS+oy$+5Su$$pKSFsV=fG^KnDm%gCRj{fmrD0wNt?j{r(z@_Yxi=5=(>LzT#Q3OQ;M*{_=}!Q%I}CX}Q*9 zDk3dq(GZO!y;L5A)rUYvH+nDMj>K;{i{cKW@nK__T(6;i*}M#@QcuctH@+?#XIDb6 zMJ@wnC8!7p1R&9VU$CO85bqYc+l1$p`z*~^owQ+PmSV&klmc!=JvIA|PKc(NY_ua~ zmu#NR1lxi%Vl8)hfm{QJGrV%mO4&Q*h14pft5sr?%I9LKFnZJ^<^!3dgco?x2QZpV z*F&dSYZj_2Y@<*gx6LYezBx>{+o)+Y1+Gs=u4ZehK}zA!F1t+AFg1Qe7b$s(*(8~E zGgp_xu;46orcsNH-3LQ_uqvOD9vDo*SmgfLkmsIs)?(>;Tb8cm5kL(N%F&1?q~znj z_{|->s}5|baeV^}QE$xAb{Wa*Vz6P=8OVa5ooj59bXBKj*Z0QuLV;VCH*$a8RrTrw zmb$>Xmk-$pv@3!H`0Zs0V|O}A-+5c=$DFUZ$mRAP3Wxt(cPgu#4;Ri*SEF=FG`DtJ zWqu-$8rsn_$xG`cgx5kDL&~-or_I#lcb!wgN&LH^VeWzR6;p#ZQcp`ZqtG%pG#~}+ zhhyF@!2yGLzt^Gyyr*3<-Z%5Bp^N4l?3^xkYbR+|@EhSV1Q=#?h;{$H9J3i6At{&;r1Y}Q}E>#W1XiEYxOMb)iYKR>uoW2knf|YMAgb? z4cI=^v=Es}-MP*WLcl)|Id?mE^14#(O6kzaM!bFQ=1DnG=GlMt-QbpQbI-Wq8oTDt;j5VH&o-I21I+90+r*6$UIVS^ zdqvkaUw5V(w4$}E>e(Fd-3prAB=-e#cbW0|fzZrbd7o0{tKEXK9O_1()Yoa*Gv6Jr z23^n`y66O0WA*%D@I2*H;Z^myz4Tg8QoI7ESy2Z8*o?qlBxHkyEc4UFt31ytk>giVRIjF{S-v5 zjr9GW04{ij;$r63U#8%rL+7~t1GBSwNw=PX^5#@ z-IsJ6EI^2Aojf1q3+2Mzu*!K^K7;2ke0m)EjybzF-lU|=-!0%ozk}ohA#oE6&+i7& zwtY=tx;NRygXBK(rZE4ktzl`yS8vJ;b}y0$i7tw$R*YFQXI|!-AJzz$&~&xHuzU1q z3D&f5HQddk`slUY;*{Av`JndYgU%de%yBvalht?hE0}tJJyD1 zJ)cP`w@_4(OROfh>l>=1@e!|!a9KpcKuZ(MJ0lY=9?51*BS?X9fX?Ov6XN&`^6U8! zqLkQ@ncO20;3Osj3=C1s6BiF_~f`pD)QoDHxpl_r|>c?W?4+z`vLLoaych#ix z93L>5o12+Ifdf`7WA1$G;cXv6A!oElM6#|%Y&lZ#!94_$vcLGtGK#0$TVJ5skWOqRK6}py=?tUy#lc$3%86iS`{T?bhjA8Cbu=gpO$cRqja5S2J ze>ndu#I4)Xw{x_lVC|A8X(l44{s$JM9%x#kdA={GLJ75LF@lJY4lkNL%DpLe1Ye0< zOcsdYCEa-Fu?b#g8}*o@vjtN|tDZ8WUI=mh9?Ve+cg8eMlZ{&2Z*{-QJ$-B+j1k<* zDlFC#DKS5|^gpyfRZ+5IX8^1VbsQf2PwUI!`R>>KaBG$*cU*9w>lC1 zF0)`;?;c1tx3rQJFrCgftoM>yeJ7_@fyA9$T+NK{rXVBK;O~g#PiH+UC4S&9oYk+m zVN60uZ65|Eoegp$#3#hp5=P#e2=J&J9d}0<{R(-7Hfz)T#lVe|o! zPZYz5G~yY(2R(eRDvj=zn={mbLa^jD5G1n7 ziHZeX*Z!58ylE-`FHeC;!DT(mH|;~rxINvpj4Fb3H0_?g#fU-MET96QCNgRQrVTus z;`XhK3`gXn=Cc(u?xP&--ZsBC3|fMIH#1BbkPHzZN?8+O6TjX$03sfCM5|Rz@Yim4?SeJ2f7b2ZbFz*&b|2PT zlhH{3W(p4)dCOM1-x$ZlPZ4(j;W3;xfi;Kok;8oki$`AxV}C;i`n& z+T#$|ccb>a$x?TSwZfky!>=DP)d08*E;D^Y6T)nWPu+|fjVEvi$$$PH^6ffwo^56) z#6z#BhD^BY@2HJe)H+5=E+KZ97#b!zUJK*qwhO~xyf z$~QW3g|A^-Q!sTN+6nKvI)jq0FQvv`!@E_bL7Vh=r}-a76^1#S1raK-44%Kf^YZ?i zPJ)9{T=MeX0I+%=Mg6mJ_OCME+QQn-{*M){AfC^v_nn_yjr0g$i-ccE*jscG216%*_P# z!8MLRh#{lkGr7y|w))CP+PaSIWj|=Hz!#luyg;YkLBzNW-)#lEHR!%?q9fu~ysAkU zh^uK|5gTKm>6%bHX z(#_F7d#B2nO~Q^+XHE9Kr3HnkPdyLJ_uWOOIqYp+G_wCA!x6Q*jI;B7o~YhOiT<3gpIoiK=r@0ow|*)vOyrY>}4WcgkD9B!g#vOQgEXVry#}Kfi5D{6_RQd04Gp6|6fkOUcnPk&^etvS&-d_Q20mEfUHh0z{8RV8_Z|Ny?K?xJ z%jdyWX3yZ8#D;}yt*vv=L3`$EI8IqOgl^NwE(^Nz>#^)oeeumTI99d#DqByZ@LBUL zZS9giQTNLxu1VPzLh7^c=-yBx4b%k@2z$$iP4F_-Ib>Pbv7;MBn8I{knVc`juNtrL z1mD~a%YSA`{dT!sv7FJdlO_CL45=VoUsN7dIzl0_w`+z!r#(+qsyf$1+R%M?32}99ri!8BMwRyj= zJ*R%-c5`M*6)MO}^V*N#XJIP!aJ=%(PTPm~J)-}%1#4#+YL)k?(0}hhBffJN4XpKL z?W}Fy`-9eYhJT%I`9I3yyPbH%)xMwZKoN2#ekHidT9r1XP)ZZ`VV8pp#Yadt>}o=+ zMP3PNz^1Yo{{g@ggy8KP-#}6lGIODYxAnZ|QYY{F9b(WZ5SM6`h#cS}RO-NMhD@LKXzau|)YJdj<68ecNgF+nzAi9_+AT2u+y^}7xMcUvut zflQt>P@2}3L+EZq0)dKd+mR1LOmx)Ff(?A)wE{`-q}BlPI>UUvPH1zu13X&b!Ye$q zpTvhs8JcN4i>20S)Wek4Ta}sDFw4u16<_>N+Ixxd zsxP!CYukU=P;YQnl|mfnD3qvF?RrET9@xVtk|fSAo>`i@cq(q~(wh*&#;g zi2W0yFu05mTg$gz9lN5I;pRzuy81%`_J|{%C~_lkYx%ka)`8+MEvZ0xy7{@L)Eu**0Whd@J-sm2YZv0ppsC zrHb@(>xe@tVo#T%n`Y)$+>Rsn++Cgo#Wrx*wQpi(24q9`^H4xRu+7o$5?z)zG`&zk z;~*O^>3K<_e0GGjJ>=7IkPYm7Kt(_*1SJxb;6pIYCm*4Z7xgR3;K5ere|Y$A>nsUVvCy2*Q6Kb>^i*8Wre{&|;8nAo zm=PX!)TE(C)BjO`KZk$C*kUrb}r1`xNtq3ss;UIwyknUF;+rwet zS2JePIgb7@w##9f`fjh>2Oi;LrQiyHETXqyC6YzYG3;YU7va&vP8Vo#FUh z_&+J^zeE85h~IaQ{%1n_?>N8nntvfF!~FYC{1?mlca-0$dB0Fr5&uB>ov8Oa!0)_> zUjSL}Z0mP*ez7Ee7yW$?`j@EX`-=5b^!FX<-w}S_(EWv=hw}%*zmu)Mr&+(F{JwMd z3#AM950rl;_&>Mv{>bltNBF(V@e5&r;_nE*dLF+c{9ac4LMWp7JA(eNBI9?6-!tuB z5R0^bK>Yvm?%x4_PhNik&eQz?`0v^MA5z)h0e?>je*wqt`_FXh7Zd=% uHUj{_eMAcw9Zy(nENVgn@fKoU9%hTf6hY=|O7dM`?-2}F7?7DT%A z5}Ha6MS;*EcWuh=t@pn7pYQwr`SLvHIYDLZ?7ilibIdWu%xeubMQSQ0Dhdh;YNh-4 z;1m?cTPY}xF`haB{^b1S&+6cRN1Wk`cPR4OSZ2Tv$IWl6-lm|)4?VN}kP`fS+Tp&g zGX(|BXXw8p4fa`&C@9G3O80JSyBp07pL~o!e5(JGG{Rin_k!;tD~4gF=UE^%sjK57 z)$DXjKQZfpJS!_bGgaUUA5ZnES1aMiPYJH9jMNORsj0k5y=C-Gd5N4$LWcX?xv$3d zwDQK3K}@2-QE*-G`YnxG{(JxJliQZi2L>;?JO7WqyLCgtbHzfI)U7osI3?*ir$d;s zku__LM|UcVdMrI6G@PcTEv>G$>fsCU9$NYuU-!LMI;MDl9cg#P#(kPzlv*o6PNu6M@IFL$Yx zQDgH6_}vTXj1(00SMWpmY&)BzuC2vdKV7Qe4qx*N?JU3{27cs zR4>foKH(mIemR+2TjKajq1d{wZI#G^(9y%b!i6SChW46%en}mZ`_#JWVw-1vQcLQjsPM_|$TDxtc zAY?ew(87OdS%>mbU!GyBKG+8yegh|`xV|G4Pmbvj)|QYQ^pA33`N*5W%#v-Zt|Pi8 zlqgT#`WDf48B>0u8qDae&%e!D--HR zlD(ed?6d39p1-@ZMfkcj?sK!(-6%|d{g=M+-L2JmdL|~%Uc8`;oPSDRbX7ZNd%bTk z<8&qvd1`N{p-)TvInrV7LjTAxa>{n9NN6y3 zbhHWAV}vvd)=g0#=|krS23CE#y0~0KRbIO((_TE$5xz#mJ^R{(!w>Mu;`+$AYlS-g zaip3H*2Qt@S^B2ACas&Xy;D~7EK?C{9#cMoy#_wZt&$}zn=T&Cv64u4wq>1V+o+qp zTe|X5B-ye^{Kw7p-o=zPFFoVu2p(PbQ9gs+xnKHVNsrpz%@IAy3{v$?cA}|H?bpp% zkJ`4)_+_f9XMtNDW+Ug2KBYBnRf`Fv58DH4$0#Taq+p^H6oIc@{v@n@ldh}z4a1*78?G}L z*qR!C&}Z+`bJSwNy;-elFIJ8bDv6;)6^+y1N1hIA(CaWoQ=;mJTnV-%?aSWi&?tR& z4t&)&=Pp&rof-U=DDvrp9X+CR{&?>10$AYtf|`>^j~LVfoscZ9y>GLc-FN+AME?{Y zJ}7@O<3YJoo=1ymtJ27n9eqIWqF3HprLEp$X+)7fN3`M5$c^X zljS}`X57UtdVQQKsoykM9%28t^KQ@i%Ei&faK#9AB(kBxK`6_lKCmRFA(Xwe`f*VX z0vq2q$FP3EI$h;<@6;p1x~cq{ug@5kqpOVdfA3H=@gJiyUYV-);pLP!dZ+cFsF7cp zMN>qNQrjr(Q4|}}BYMqo`Q~U9e(Mh1?#LwXy6gv?WxpDaZP6*Oh4JbZt@SNS-tC*ZN1_$%Tmn!Ygd(%?kfLiC*DaSE*I+%j0C0K7K*A{mQkZ_ z;D)g~flJFE`x*SRx0P|lCiNBWrO2^Jgh#sO2o>&GXf3>~|C@8oHeN>_`6K-D;`}E{ zc}<;q-4Uwtt2OJNRy&taDW%0rS!3JRO(+u!zwtZf!ZXMEbbfE=zvpRMo-@%UD4p@$ zUZ(eF_-WAo`qMa^JR8Z28L?mwTd&uAWAtJmu|3;8+X|nXeK-ikY_=VQx{F}zui#J9WAMWg? zXJ3)-o8{w|9<{2dvMr8*yBoby_DN1^FGqS!*a{afuxy&QNuKoAZGY|DYI^V9i+S`& z`%C4p(!d_gWl%EO6hK9+Lq9N_(`RU>L`hB13wW5H*Ok{C3C!x`vF}m-Q5cPjGMeHK zTcVL_udIBOTvf?tG}m;6xF%=J8|B8(+Ji+AjEn?`L_Y5T1#ACM#48=+O3uY)!7{_H z&B0~gBvSDO0S%n7@$VZyf~b63J@k)|=Y&4=8t4$X&s=!d%c9eMOtCC-YG7kL$shUK z%PFsAv*HYcv4{Dm9YjZxB#N?B6P%dVFjDdOjxWQfO&7|=J66@UJ7>$gmwUnPrKIiU z$3}kVC85>aXY~~pZPE$f_(E{>n(0v_l8(7I{mu5-@~PylE|ZIkZ6TC?!Qv>&MG5yy zOOxVw7XB7JN^}&GS;FR8V|{+jP3$-Rbq~Gz01}B_U=6qU4)I~UI@BdO#%1mP8&a_z zr{YK~$=Aa_LA0R8h`8Lm<(*qFi?2!Po#nRn%0PHe?nqhFu#kpoy1^akq+m%dx*)T> zEx3vNxcDyJfzetv1t+k$e0}tS64ANfrqRgaLs?vTzVGOpVYT>D*0QmN{2IBC!Ds-w z+9tyznyY9)z3kWFZ^6;4Fz6crbQ4Jr6@j* z-ClK1YSL1dO&2w9vjCWbD~wA$8CU;?fna3cmzVT4nAzNG#{^(Q=rz`UeF9dPZ`v>j z%X{|e^JiPi{ym)Y?c4oaz2e#ThyM(J9c^irD`ee^i}PrTUjeUYEq zU#iZS7hxWc$kO;&#pVCvMSqFCo{+@J{lD7{>PogWv&DN8b4z{Yx>}@CT13oR{vkAj z6F5*}FiCI95O3U4ZWchxj&yP=@Hn&oYw2NlJl?7&2R%}8Ej>r95E)1>GcMO&KkmUf z=-*%IhG!DDdNlGW-GuE;$c@KmCp&~=6&oI~u|D?@kedAd-~I3lKAGQb(|(qkx(7Un zR{hSG)XX*m^^=FU;Mmk@vohV`IT>e8cmBMo*YKd;K(RH7@Nu`7(EC}cf3)e}?kTX# zuKn=|fv4cY`O~p?_b^lMHvxn=6JSk|&KTJw6;iiJ0hpm|gm1Z>k)x&=W ze^D^Eba)RulswHshkr%Et)Pn(Ib0&~Eh3r{=bcc1dvW*xuZKP3Z?1aZ8{ZuQ);Ro3 z4o7W1KNdRtxSZ4NO$9|2NxaRlQwLAv*6FghQRvgdY)*p(HmJwpBf6eCYd*z_qxMpk zodwSi?s`mwZ=P^XlI{L?{^hil&MKar0y(48*}8fa#2`jN>BEmD^t{t3T5$9{)~|gi zvoxBeGbBxktC-kT^vomoUt7bYkX|m({H>{D+cuN*t=qw;YAK5*5hl*zTFn$NogPO0 z*pSe~#rv;C18-XjHJe0;l7_#YXWAkesp|R}aGw`xc{M(THe>Ru-^J7%$Y>s{h zAt+Mt@Lp-O!tC)ppxZ&bR_x;a$i11LwjjV{nn+{AB zi$N#Qa?KX|pCoN39@ArHG>2Dy?eTW;Th#L;4ZfW168nOgI1J@FPM%iuK+db{veI6) zTmm<evJ=qoh5R%LUS5IO(9$&P5xPxV<(u)O}+rlm^UQcVNdGkG>IfUou zy+Fo%;)M&iserz|GORULUezxm+W7F+8B&X)(XX?0_rE$@UyV=B zzCVJ}{{BMq-KKshlasc{Y=esH16?FgU$U;Sl zDq59k_niT z$q5Ie!`~=}S?iYA8CMy_dfg)HI%w1?aDsh73;G2wq7>U{u=u=px4G*556{koo~2-O z^g>|A`7RF9hY zA(-Bb}p6djMR)VLVW!!+-y&RUYcO9%)tIpQo$3B&ORI4tM91h11GKNM* z`i_+HANC;RPsva|)6uJ3#m?Z-9vo!yj<~Yi;%nh-!)n+_MWsGpS^>+z4^l{BPdB>YDg(s*ix73yjohR5!3 zceVy@3i9KGY1r1I?=M|eyn_CzTztJLmn zm*H>@Hb-tMDvIuKTV@~4>1MyUx`^5)FKBF*_=9c169!$8y5D!Y#(Q?9BBHDQFLTqm z$bT(rg|UsxTo1KFed?p-Q5)DVGeOPfyq(_3tRe=n&WuEn|rVNJ#lVsM9x;&)33GN=T!BB#RZhc9g)WK z2k`y2H|Be=lCX4Y-lj9WbmPH$BYk&`_Vyd8#bJ(I#nuL`sI|4cbQIn;>!#{m6)N*X zb}US2Gmie+s%uZ75M^Ke**iw!x zmsjagOIS7%+(hKbhQ8XrGUY9Ptco@jtfHBEp3AuCPg!{`PAV`BWE5;y54W6xaU4Xe zgKoO*x{>2zF#Kl`cO+x+0LC?0@C46q@=dL}+{MN;M@bN_wu*m?|7p1E>DlX>dp1-& z+t7lZeV9;~VN`oldnX?C7W;Cl z@5e@)$(u0~MwwAc$|Ds%e>@*99N97dczcH_>H)kmx_Dc>#+X%B8dp>C@k1{MI`sn7L9X=|vitwO|cV66J99N${VR>l5$pHxlp3 z%E2pXA`dZRKx2Qc=4&l>)bA=k9^0<7Ag`e8FV~)H=UETCv9#cXH5jS}t4c-1zO(ne zk+)Slg0*B;{$4JDI?1^BZ^1dyfIIN*_#S2?Lh=QH;3;fm%bpMO%L@T zrTJOJQt$Av)pRTKC$TJfd5aRaDo1;w@JIet+`^afhe0#-VbkNiB-Z7E|iw!)|{P06Qd&Hdx z+6JXSpWKvihU)COSPrmTe||^1^%3Z0?2@-%%B8owtuVuk*K?WALg&g???p7do^J0u z-_%|4vbVwF~uW2RSI4(B<5l~Novc~2 z!xnDs;35iNf_D}|2-Nd;f&;)fUooz>vnQtG-jP-wr07}c4Yy-%-8?)_YX28SCDXmB zEzFGU3ZwO{L=`$2n@x?Mcv(!n%zlh$HPll~kL>Bg)Uh9g)04vguAJ!YrWqj)>xopT zymZ_2JNs3u!LR2_qRc$*VArRAV66!y4@zqH8Rc6@3WLS(MvBndiym_|LeU^VOQu~nEje*+Z~)Wxx#G-Zt=U`` z8Ohd@zy3G^hBzwYV%|6FIT8PI_HU*~@Uc$^cd9`%aixI|J)juqGjGM&?x(7o`M%}8 z5X#NNjV(7dz`SuWernvPu$<9>t(HyY`?52QMQ6Jror_ct2L3j*(r!T;ao+wEVv*|- z7o`x}LGGBi-y3%qX7u}|i4id~GUcYWD?ErX%wXrvc-NMZzf6rz%`x^Ueg8qTFCKH( z<1u7`>shbg&{{73c@zN~Vauz?v!$>0(TY7`q}hR~OFuY>sVv`Q3hpY&$)uH)zT|hA zr{k3-4_z&VxlBk;J5tK+rg?jga4?KM5dA3w3ej{fXTVC+Hjk}zxCOkCb4|Di!CCEN z8mmy0cf}3GeW>zUI0ra1nAzOL)tExx5?CsE$9e-6$}P;wYB7+=g+LWuhdKhxvtts$24xVoc8nl5+w_i9$x9*q1JJ!24-cgpIl%(tp%cwE(Gk_47 zdgXGve&J)Qzp@w<#l8;@TW5i;KE8hiv6zg$qDb}tqHPntBk zmAbX1+Ff*cpCuYT6&@XzyC|8|{e1L=^~aJfuHwpqkXooyf>4NORFx_f-C>&y%bx4E z;JoUqpN}`pjpSV9mEL^;wz}%rc-Jo>O1Tmp-@D0g^(Ri4ec5Okc4|UKw+?y3;|6YU zMa3i6x_%v-o$l~;=`M@MD?M#KXnvo-xrqHnoUNX>^Ut5;cj=jGYP>JbQO>Y=NqmAp zK1V56qjS^eYm<0tIHTiDxdouyHP;L~+)=XB<%(j`UhTN+V>M&^oR|Z8m58G;3_N_Di|@C0C9B`nZDoAJ zwmL*eWs^9E9QZ zxVbM{Fr(|5pcrTlULB3=ZY#9BN;$Fb>El|RO=<4p`8tuKUus5x;5w8 z6K)KKT^R)<<=J5XX_vlr8;)k;ZSVf|t|ZX24ivevb(Nln7pNyR&zgD)*@8;VNMN81k1 z?%$CWn`i2E(F2{dkrC04;EmA99bKw7F-eEeFN$f$9p>V@Hr!!H5@{xo9!;t zwkLP=PwpN|KSVp67hQI=2~7p#^7g#aGeJt6rB=b?e#39}4y^>y6h0OCHa;Cf4gs~) zec^_}0F~=C%D2T8Fw*3qqBib_N_|sOF14L~kf+ z!_8#zI4w=C*6M4tIb5S3$r|6Ei;nlWzEDbvL4?b$w;JhNMkmMl+37}D)p!>9vC2xM zhJa%x&a2^hnC^o*-GmU^ilzL2w*2_PF^iZqdb@>qYflcKm+`Hsk{)>gpF-Q#7L`2I zBM9f(#(6-P<5;!z%;??OnWn*AaGP=mw;6c*>j)x{44ta2Sf~(XkN=$DPPw_p`;D); zV#o|WRm>-&Mvi5(l%5`y`YoqUx~M+fRsGV4rQ;SW^YM5(9%-1Nbc9Zm!E%z=i5L3yQa$cw6Sbf%drHN_~ur_o5PC*xdO{Gj->`e zvaL^?2S3OeskaZ}X=B(s!eYfDVD#HEkGBv#E1RTuvm-qHm|Fu{IWSSIfHjZ=Jeq}M zWDJ1F-}~>m|Ee3%$py5T@lQ<+UHVN>gd`tsGz z&x39v^_cLKm;^f;WBJ7)x(sXr;u^1pYgA4rieT9LAm*2DW@7k@*pK>*K zXDWRqvXf8&MEdP-1{*a0aHGgD5FuV6oW$Vzlcd2<&j;y~KyVND$aCnSN}|wr^=_>7 zQCx?|UGKu(@!x(aedWqM;EBku_O01}X78J%%8ivolMzEF)1qo@-zm*k>CtL^Kh`xN zIj&Uxvin-Lvy=lLVsYoO&eay#)KMC)_-?xgZvsmS<|;~}CMMeQr+JL`q!ndKYr=%S zeKlI&#Vo%~urP>G^y{iqFHW~4H`c>n_lNgTnPW-m)h{E}8*buz#B$xwXm|KcY4c{U z#Za!s7~~>fx)^gt3l&M_XO&v!xCQAQ%Z{r6;mXD!(hU(aC`7gDDtuaFQ{1ZiTb1j> zrh7nlx!BDAc?l-gT)tHG_xv55^nrRcp50gfQRMi~Dn$Ntn`FSi>&(cC*)!>R8% zWD-`x0Ztt^s*F-f?2E&kZ5{U1d*in4D2a^uevxpYW$=;_SNUTZ21vOGQjTkY*OrKg ziD@%S+KJgK0hR6pda8NvmdgIQk$oDbE25Nt{>09^Ys;r9VMk1N`N;`f*L@yFM@FY- zGphk4_=&6$Ta0B$1)#0ZN7h;jbM08uGdLnXm`zr%U1mAR#Kgn~udgpu={j#E{d?4t z;Hh(RSASr`X)o}E?x~qoz@zYdXKIPC-$Wc2VbvEk3gzg(?vnh+o zNiC2y99jpQV3ef^#aT0W?vp(|pM9+PivmNHj_AOKS6iB2bG{AvD06|;dHSE0G}CLL z;0LJ_%Q}x+-u=5DdRcmJl(&`8i%QST(d#Z0>-CpB^aRQ7@2QdpTb8ZEPI`{K8~!Qj znh0ZRON6HF?D5zyf82$@Leu}mh=dMp?mM73K?!3Ftn5}6yejrY=gy5?=f zM^%{O+F&u%eafeLVO`f>cQ|l_iGa;WX4Av#oDoeGIA{wZy(|;V_8@KY8hbfcl1TPA zGCL1+7rCjme*Epx&sAotkcwYs8rd(CW^yc&?%xNMRPehNxi9yn5!1k}BxcaL)}FN3TQnu%2?#=bALx5DhU^2gHE)gRs~Gou%ZcekrDYJJ|cNv5h5 z6aFS9#d~uoGxg7MRHsdrQW9$Y&m)Ma-7UNblMdmE6I;7HS}*?@kHfgx*acRi;Id+B!db z+O-|oZ~P%zT3XV7X44g619V(EYQMm}I7A~S-wDFD&8q(9BC9g~um+{xr$=zL0Jo@9 z?k>HSbj43L_JKZP^=H+wO|)N?(!8exO$Xq&E7G;Zp&pm+mgYgNX-dRR55u+>QUb~A zH0g_*faNe^%|rd>aGsEUkZ7>O4={+X+vg&h^^~&pxAWTt&v2CwYP?mC?HB`bfoXb( zfBZ@3tNkLVVfH9#2#;ZG=UaH*0UIs(4%W&I;GKRtv&6iFF}Q0L+kHcDfIh9EHneD% zx^V0WIhrl_8L9Kz2`JuXwSPKaZr;o*`@7Xh)rd7WE#DF}zRr96lhse^cNLXh)IGWPbx+V8C|I|$ zqmtZVjc2$4rF#jsH6a}_=piwmURNm>K_*>o&hl<>{?pGRD<=8DnlB8~p9qtq$g1b} z@>)`8aroS_c*9CvWOaP?L?m3ap!JpZjk>#%Zn&H^J3EXuK^yr5C;>Y7J_Jz?@( zetC`BMpJVkPsr?e-Na%1(Zf-G@2+D7^f#Y6buKMlnSppWdvC+1ZoUGc3Wi|05_YA| zFYL54ak^W!ZuJ6Lj^f6$-z%f%?P6t_PlAY1I-rNH7k87iZ26l(CP|;eY)9$Z^)NM( zD;VhA+ozUndK)lp=AWOsSBqVR;e}px;x$k2cP4D?8j9-F*_k2KM^&ln7!NROB786- zF3`&y6F#-0IN;cz!dH&S;jjIn&Og04+Eh4k5$p&A_7@qqR#wA5gBn{a(&mKtd{I^u zSHm`sQWUh_NT7QndXQDFN3G;8)35CY!{SL*1#0r(R34SciuXQAz6%kID8e}|7omY0 zouB5-Ar|DOo}>0jwaoZ)m4LBbnJuO!tr>~aBkLP56p82d8;Cky%%VXKL@)nbBO`nB zs{yccAYm_#?$YwRZ+m;VRei*?aqMGTn%dJ$uo%&?(21}$IhtQ8BX zWW{%?T24`zAMS_d|1QB?#0n6dK*!je{pBnq>g|b6s@(J`@n{9iMRqsC(Y!#21!b1j zMo`MpiXKfBJ4{-yf_#`Ge^2bd)@N^{+4>)K8rUrWwm!xd9j6{zIq`GB@aKJ zZcCzH-~rBGLzMP67u1+e{A2}4G#{u)BU>{clD8kG_>Qoe!YfZi`K*+(iSGfszN%+R zMUpTP-WB8)|v{%TE>oZA)oBzVCMz2db}k`7;r#uk!&c z=X$W5NffK!okk#m40)wR*0?RYLewU^BWqg<~>dwRDVC394xmpyDd>_ zH0L^n{M~)thUdbK8D5yK6a*o7NimjCcPD)7HXsm3znM=U%WIU98q_OISH7&=SKZur z^(mA)p}S2_b@kJ$6P{)Ew&K_y=5Pu$^+DkL04FZH|Mhc7n&IUpV8RvCgs7^B=QoVW z1w%QK!r9LE+nYM7ax+P*n<_pPe0l+-B3>gHVnp@p}+XQbg(=^os#X zTP_(Uqf8l~0pjDn(S`~4nvVu7w;;eDc|*h!iEQUr%8lJ?32M{3__f&t;}CwKTK))p zW&$8oFd)-Bs^l;$zBpAG!TscrZQvd}h7 z7Pwq!8h&IY3r3TyzXfvL5NaKPpLXrY!iSD+HGK$IJ03Fjy#4!X3laP%A@^+o%l z*3Z&V-sNdE(7f1eqQVzb8*Bo$y$o-B!L&UjXKU|uj+=%rAO)8tLf^{Y-O?czJ5nR8 z@?0MonVC>5eEyclz5R_Tvxw`c1frrNr^;@E$mPBSZ70XINcBZ8&<-V6Ne@4QPkmUn z2p`dNy06|klo|CSpR4!`b{3moZ8h6t&75QFA{2GxAma$#acw!r+3Kqa$%C6AoM4j| zybGUwZI4lJ_qrlU3&>Yh&C{M}-$`zC=XJ5EBc4()}Sxzhv=I zAgJi#r#txj#J=xTTW5n*9dYzuyl5M*?dLV9ac}f$QffBN-me^9>bq&&P1U{hT!Om1 zGXwD?0&nn+Cw@L8c3-$zc5&`Vy+BXu4SraGwRY*r8?#lQ7~4L;SLRBiCHEZ{M^iC3 zuJTJS#6rFIL%>DRQ;uxp%iO_>ST;&UM8324yOEc#wK`iF%uN>)ZmEE|WiM$+DA|*3 zfLfaqUb~!Fc%p9O0DRbpT<+!BLs`QSO;Q)^I)~0Qq*^U8Z>m%7_~gjwuSkMrSt2qZ z-T9;aCvi8TWxBap!ntg#)J(e>ZAeDf;m}7#gYJ&;$dg0*b;4ls$2=@iMD4Wk5FTnq z`O4il8Z#GnsMzF;jpp&9mW@<$Vei-cZocQvP_&xK`zKR6Y5xiL!DFAN6i(*QRgPzwl?SW%q%51SO{@Zn+ z1LrPe25k0rzTsfb@=ALUar9Qle(c@45s~y*q45kC4@3+TWEucIn$Ob&SZkcHcnPRB zu|?y!Q3gE2B-?jukN``czXlA3mgFtOxOKY|*1fps@i=qvs>5vfWZ(Pf1n+ z_7s#PHnlI$I%~_bSNn6*O)koRr53afQ26NA7K&lXVRRvcc|=2*5o}Q$h}p3 zzRD%Y2;4I9y9a)Z>fJImu17$^CVWN!KS^h08uQF578ndZ#DLxXDZif!&>CNp=>y5` zLE*mcdb-?e2>^J1xFfg`?5_88&6l&?%tu8}MC}R&QV?Cp0C54k2WJL4?i}xzc0HM2}L$92ig{FR^jELfAa=H|cYeOi>R2{Pu6Xx8&ZB1(-KuTTewq z2W53-Z@SlO1}8AR^CzhD^Gi1beJcsKQHh!LT+2{k8q5i=k}0#H)DT&p*%PJEw|ZZn z0{zHUNkQmo)Q2xg zW<`#=qEsjxE*xw!)dVJ!2R(D8UhN-Dw{ZqUI8NN69F|k90R*u11A1_m5E;6@7Pjg8 z>ok48e>P-4So}y@34Q9)lcDwm)Lt&YVsy;*U`ab`qt?XU?fJ1oIWkft8(XVA_vJQB zclUK>Pa?46ZBgzV@CzJk(+yQU9VW&1Ou5gU&X#Vo~4D}N+sns z00YE)tzJypOWXck33^A1PoUqZ(4~x^;8Fei{Q2UCpHH+VNu$6l3|BGMu(D*6M9hp9 z)Q|s{7mqs%V$5#Vs=pq+nbDFats3ujY06ls<~iFQ1>`dES44DTQiqbK3|=c!qk#Ii9O? z*+$i#*?PrRT!MuC!UTwF3>S~bA3~syb5^Is?gO|qzV|Nu%GlWr8kVc5+Hkq^Hou-8 zm)-b$(`T#M5&dW)R(lo*-}VOVIo1TE1h~>=QBf^IiEUvNTH19^!_{@y(1Rb9hg%q~ z#Hy#<%+XAHsnmROpH1bt&C7oiG56v&%#c*BG$)g=qCF;S{LOyc8F0sa7|tnRjv6uY z)Ug-5fw^lASLN%$kLN0G^>hF!93l9v3qt`rhGi*)iC_f^0nsVbtd(yf?=nA#KpFcK zL!RQ!r2`KR?AGARMO(CPQ~25eX%4YWtI6An>XBZ#{3?1*B8ARkwx3VS3d_QoR~r(v zm;N~Pv?k@u`k$kIjAn)6To5-OgVZ(KqvIz}gZV22$}k-aDPh%;q__lA_8;Jw2x!#* zhsv&(FNuU~6RXi$T=~QrvTkoaqnHL-Cx+;i`Q{*8V=4*;FVGe0n(G(jfsY z(5L8UmKwKNthSZWXLZP*AYGx2N)4?9ndiihA3x${y-J{fA0OYY96oXUc$R5H2$PKa z2PQF#2XLK(p!f&PHj?*+_ZN`&cA8{9G#;=wz!mjpd1B9Jgj45d%*-axv)_BE%CZux2j=$NvV$?D`ycDo>1Y}EkrHP7(Bo4s|-Y*0{1CLIu zztlcuvNa)-MC{|gbm=1|92m<7OKh=V1dmrhpapZCVL$2|Qs*$DFr#a?9T$VN2Pba; z#A@Xh3<+YRzO;??(sdY4B7h2>R+BdYgI5L@DNyAr#KIf&PEZK->|`#{ z9>W`c)Tdp`%?;S>hUfluHR7m)>q%O+?)UO>c5QR{~@icXk(GNHfY_h3{};@ zGY7d=KC)tg8B`-2BtQdhRf#k7B`f{4$=apbAL76qV$vBd>C>eo3mS!9g-QIKKMu|v z!K^v6QPxaZhO$#SlqbjVtT(-N*^RCN&u}U<-@T8t?a`=&bHMnbG@}cd7Gy=OYfD(v#Iot;e&tD5BKv9sy0r>t zeYq&dTGSG+znf@|A!erd!XG~dN*C_=PHsVq6r>5U9|BI`3jvZZq|W*|#fYrtNmvRo zpf*;@lt2EOY_4q1<$6p5l-{&IS{y|d4U7@})K;3B=k4}=2S>(Q!o$OhoF-s`x&`+X zLs)?QQb#8i9HfjATS8M`UmwQdz|wa;6ek)GDMqB+6M;(w#j=y!0yH7;4k^m@T@4t{ zt0G*AkWsF^k44KwFIp*G!sX}+#3K&p=Tr@q(;ZJ7juSO6@+4W!4;Cv2(6EYR8-n3g z@kihODz+ar0w;9de?G^b?Jq{F)RS( zz}g5L#-YmUn;zT@M{Xt?6N-Vmv&lVnws&EB;94I0hV`LvLGSMhpKT0D)Dj=>lsdL@( zhaNqfdHB#9S%#SkSH#Mo3w2Na5r)1eg>OPb`v9B}_uKWJA5HOleolF~!bSW+9fixn zknYEiAGzcm57O?O&0b4sf40SELurAqL12<)daO}5zu4@zlHi2zY66}3PK~?lt`DYW z)cp?7`&)X%1@43FZuuoDO66#sE+-e46?ks((r$T|IXG0UP?fH>1Xkh7$20G#==oyl zV)jP}o;V4*3`+tk8?vCO>Fu&|4O^GK1^g%7waL34VsoHB|I0-R@$Zoz>>a}UvGW

Z4mf9J+1u}l-COZh_xLks7FTYx_uGI0(E|NpfD2@$aL^Fj*fF5FWoAC zo9|WrlkAoWNsa{}6Zi!O5}s%D$S=C>Z!8|gmA#lYQ}f5+Mu}dxfG-7U1ItiKhnNY_ z&Q%mUi*I;T#L+z|Y;0-i1_#OpJrKew9hsD5tAJ^W;>UWeEzEXhhFW!JJ%nFBvtL)x zdS7ddL5UoVx;AkR&9)7R9GQ4X?MYU^vOphtvZ`O{0~5*c@z{5-`&8?y-p8c+mF8 zb^pA)Xy?}7B$C4Brfk>d_F#7?H&fhBuPl20wfO7dZ@p3Cj@nrjFHXcAe_5LMqD8C2DhR5ijL38%x)=Z?3x)7w#f+=XDo37h-ES^HNl;riTA0P|8)@ zszJ$JWJtLcAr%^Y^6!;#+O3uOs|`_MQKs;0=g$|b9euL$Z1o4OP$Y)Ae5eK4e9w0< zb()F_o~f23<}h{(Of|M+91d8=rWe3s2aDec|H%~$-;CaX>anN+#4j(7HB2x~d%Izz z0clYh!`7>>Of&O6Zu`vqW>gi)qE-shIdUjH;=x9N;s3t5fx=?TE-p^)Df>*CQ>Qc< zLRiJ!e?Jhn_zryDf(H-IYCy?6vdZDT$tX)sQE(TyD-U@&$`MIgpyn+&wu!EE#iIA zzKWnRi^)sG^`AZ#RWN7D%l{d$qdTrtWh7ZZWtQwynFC%jx{$wKkw0YivB~b#NY%>} zph>E`JMmhaCf&lK4WOV!uc$7wIEt|E!Ox#A{YPKnzM^<#ME51t`Y zHZgM$8O=#v4{?Gom5NsM*$9!DjZt;1?2R0Ad=XMH{{at1sPG~FGakGi@%$vMRmXdI zy+X5Kpc1R_d3bt)5J#32;4Hf`xa4*BAx+&(x@~f9nOmObQWcp;dbbQvWqj1*Zmr8y z+k}_{f#lVi^HjSXcm8 z_0JWs9;qrfFw90ZU9Q1GU?vf{xylZh+g4coA6$VN>ov;-e-#Ip4Qdbg?KQuvz_A|N zy|L|jO%tL-@qMf)Fvy{lz6$zzAap>pKIN4aAng`j_gu+}7BbAC=SNPDH%4%^?hO2z zwc8-vHMTUo&n#wzotj(r3TTPgIzEstQ|Q zI>PBqN)<`-6$jV)!I@&@r}aDS%=qwx1aNK5B&f}HEM6{>5Fx*Hpp|48?thnKRG#s4 zu;*(v0bMo+kS_@BMf*d-kE?*dh)3Lgo{J<9+K;o zPyzUzP=c1iev9Ds^!ahy=fm79wJ(_o+E(yZ6q&^U1xe?3wy&?9a^ULiRR-c+m3ux7 zs}z_D0q4o;+uKV}KeD~EqZ+uhFTvKKV_2DUG39uky~z!N0gwexZk?3iE9l1Y>(Ik^P$s0M4-fo_c$> zKR@m0v18yeL3_P@9XcX{1v_4E&kH%kcb1Zr<~wJ{jk0#!#k??I7ro}3FO48FVp^191Rr#$$+kZZheLsKT9LluaaoxMu8_4O7o zX+?yY0kfJN^MShqwLlsKW;NS>c0BOG&d8mi;R6$AO6g1=jXYD`1cQkveHVTwlx!U% z27|Ex+(o>w=}QUw5!{zA+w)z}F&-?+eo4!2M|jEj_op0<#C_ubiOIt~Cl}XM&*lb03O?n0`I2E5@Oge7eMy{84MqL2M^ug8Hns$xATy#dHC48pQw4 zCK%Tdjc~POnQTxQL1OZtqNN2x7cL5ME(|T_I25uxmM52f|GvgiXS|O^$Nrj;j^W#_ zE#(8pAj-ugMalW9)KP4c+iy9#l`+~SE7=xITRq=b+_M(GCfuhMzBVeVKhAx>`T1vp zfo-8wQ58Wm_KBIt5_Aq7Aa$(#3YFUj+JNPPDAQqTQK)CGtKV0GHLd_xMC9)QS=jcf zFq&=PA-u7kTLa*LtT9D~ZOQeykDsKi81U*wDxf-vf8Hc%-k-btFz z$ z(tyFk<6Gv`Fi0l1Muxw}e)mXJYWj{FRht1G@{-Rxk}R(zIb1*1EI`EZd{sfiX5Wt% zk-F34OuEs5V^ z!O1Tus0#65tP8aOlhij%#SG=;Y|-(_mAeI?S-y)G>w+jk_7#9J zb#z0D#!^)u9hlGp0~*bsVaq+iYcVIbLA@0NTqf~FMiSxR#39Xnd9Rx!{PIxNbDa5dm8qX{8o z=I2ing+sk#%tn_H(CPVmp5CT*r{^qQ$;})^5l6pz^8jz#H(;g(+}XTh_zEydZEp2> zkpQ%HxiSf!bZgGmjJ>SY^k;o$;$#Q{QTePC1a#;7_wS29#oF0kZkKWze+2>>GzHTA zwZ^r+CC_)CHt#2hxstvE^+nPg#TSXpnOJ`2x(y^yXr* z_!>PxA#1nBw-wGaPtfkOPwBRadHs?Bd>%jUp}CVKfX6I1e;eW_h;L`VD&5(6u1<$y(vb4>Udsc4i4gsiDJMoC`XE>Oqf&DgHG3|ll__;9|i0$xL zHX|fmE43MT59mq+E;%{b3V^tI2lDDXfuHA9;fWtI2iNBD>1ziLEbuh|y-r_1JTfe5Y&gSK3APYKPOvpg$c< zweu(feSL2N7r4trXl?iW!dFvj)_~jPsE1ocBPPZe07*!-&oo=;db3PpQ9q(HF+Lh_ zD|sPrmLPVoof?nN4CvLuf4$`7jMilwQwd?0%L9+vYN{DTG&$th!~t91UTo@*@yd^A=urjY zD;$79C*zU0E!Rea%X6~HxEo%DE{ic~ zumQ(F7;)Kx02)`w05f6U*~~2{st+dOr-5?{8g(D^7SPun*l2DV8aC*zRC<-nh%|1W zytyex9j0h^qOM#A+jv39z(U{9b*^6+NY8nO`aYX0nu{aVw)nKj&B6j)6E!@61g~aL z(Kz@f2($(m?Bi_rK|te(o4u4!%P9@NP1_#6Sf@C_qyctCef>2Z5Q+v%a%`63E##9` zggIQP`L--t5Ir{7riD2@!-ohRka!YsI03>9bfI~BNs_+nWDBFHiFqy;^d%~`F1ny+ zR}F+u-DxVw7=;~SRnaNgr!>*`)jR|D2AF~Nv4f5W0JieMiJcw&4QfA_9lg+EU?>#`9I3VM4CNZoHep-EGcE%LU}fMz3WL8qFQV(tMK{r z=fME}!h)T2rp@~AA3(G7_VH;w1xs==ik|pke+fXc)S&x@JLx9eG;lDAhDo%EfPtw2 zp5J-u<9<8MPT1aWn3#a#K$|72tsD_+m$>34Tq>A;fyQ+H7hzu=)#bKs4Wg)^U;v5| z7D$71gGxz*(j_1b(w(9Rk}BOT-Q6Gv0zX8$1?le2?|E_GyU)J&eE00ZAA6i}#$fH2 zwVw6N`OG=j{%UXIt-6BVopNUE5xL&ka zv9XlzQWWDtWjV0-xuwUB)8>scC2#Z$$O9@m~mg& zUmqys{k0}i7^&*WR#3^Mr^5D!k;z9WhT`6`AMzKa{y>#v@y6=QaL#y7?E(SI%m&MR z%SY7Bwm-G89A1`JW)o?3jt2+cryaeWabt$~=X0hV#9RzMPz^cVIcN_NWmTv!>re0i ze{XMX^F=dbJ=j_48~^5058pIO0 zuPYOhWQ5KL&mjYQFcFK2 zf*8t1!2@&d!yD_Z7J$`Y_&jG<05w?Nz`y`-U?dcKC~=g$yuW-!{#`qjd`)M7>`%`c z^+M9n{qv2QEad((nZQL3InRqzUa5BNw4)j~pB0bzlF66tSTmS-LdTaLEgXpQQfqrj zkAUA_HD#ZUe-(VmPHj$StPLl7@Sy+Cq_Lp8`vucIdq>9%K_7x*E=}mD{;xVTOZ>uPKH|7z~0r>1m- zq^V>=IL@Bl_hSO!oijz+jLMX#!b5yD09jOH#XFOH4hh3XYMQc5T-O4>AeNpQ*D3{4 z8m@#T;UDk^=lT9TaZsVd5wsk}E6RpSt)sBh*6l>rR zGBy+wcL71<8}}lC#**Tc{j9y1WU7)>x??Kc85J5#Y>P_ztWv}P+C1-5UsAk6c=uYb zP~#jDUjp46V*eR`&e-Oq*0lcyaqU#OIHv-0lk`$UJ{bK69TkIEnhtSNv}#I`aytA5|Dm;FfpYGXpQ-rgbd%lL zScMjYFy$TnvH~zHV@bEBL9jlw~O!lEt$OfOX za*$5wA8DW&;s2qd`Y=J{Bh?3S60-hwr?omC|4-!= zS8*n1_{Vs}nUvUdk9ncwIjGsdn?Kll`6rk*77#CYzJ!k4x}F@N^sR=527MqBB3O)0 zc?A^e5I+${N(jaVnyx23*vZKWf&^$=@~{2j1o2WbFD26)h?pblzNJ3$LU0W#Ajb63 zT-R2XhBZCxw!@{MMk(yR^i}r2euR5O`0us%(mWQ{CPf)d>5OsAx!!p|;sR#CCZNfM z&gI`WQNgP|=}o*(l;r<}y}4<9{txq2*`~Z7GzCAX4?ypuGZV(3u)vA#kSj%t*l>Xq zR{~Em(!uMF>~(T>pH!?=7i7`-`1`Zow}D>adCBV$W|%y&82A57qUgcOq`?*mH+}Lp ze=+}0Nh6|JDk<}J+))P8vKknFK#JF=Pkze0E|GpueC!y=DcKD|K2@jhHoD=~dMIw_ z*buD|bG3$vbtl))~`3qBF?ELiC|+6BTQvucZb4v@lwM)oIskXg(k4 z+lZ?{hOV+9WV@FVXsGY%Brq3X+FYyTTD1c| zxX&|DPgOvivy0hWa&jNb#jdz=(3c@i8|kekQiSx6!Rec~Fk0zOK71t|v9bZ1_G)Cw zVrO|o8Q8(;_ShWAz@tX)o56Eu&-O37BRtpE%IK(#dU$MXETLeU4LWse!rRfLKdW;5 z?Qz4H*733YV_kd$ATWzdhKh`KJfB{Jg}Ehs=e!!W71d9yWJZ9)N^pzo$%T>MP9e~qVKY`S`X6>dhOl=a`JgpI8(785!6p|9+j6b<4ih);WST9H)3_< zf9KN9c=DO|U?Q5qQb{l5+wMsGi+~XF$L0pzsh{kZz6Q;8qRP8O)StPC_qxb^%eo`YK_SDXPm|od5(Z(+M}2`{tLk}R1JJZ!2h+#v%e{)nkE>j5?rZ*79FKU1y*-wAD1uhS(ut-s_wrF}=F9=`5! zkqXt=bP1kWYeh?HXG4;|H9NV4DLqdnM5J+MFl1;ns|0LkB!Y%JB>d|oiY3|S&x`ch zNse|W0+0z4G%?$rKomA*8ug%KVD`fzb&43MclzHekq3R~GzPb7v79c!!W7IjOwN;U zFZ4+v{yy4*kaB|R(*>FFo^q1)DnZN}_Byc~#7S0HPeP$(wCZKUdVqWndleLYN zgqexo>5*J=JqX%w;Z)VfAfx5tQb9W0^{LjadR#tzkhgH!uO|@}ZLF?J>+7dF>?~)2 zqGzs!I;Sv;@hPF8lKhvTw54$W#%m(TcmCfYk*fiSS8TpTD)?%+qvmtB3|BIH+nT3{ zNbec1@Q(=?zrpY*<$62sL8?^O6%Z65pnGYAVIhwitk{8CjUw zBW-r^QGO+rLc8F3<9cq=V{xv=WYql>55yf<2GLg3=JOL^@F*$3V!B}~wYOQV0Ea8O zfWb28Z(Y_bm;*O%xi~oJ^Ne&PJA~7r(DZEQyD5+M7jh7?8}4FV6;1h4hJs})gORn) z@{MHc^Qd*zu;AO7G}K+EO}2D4gkQ;YF5bXDsnh?H@U`-TjavWEX>YSq*kS`p-M5ri z^94N}4lX~?Pb09>5^6k`=hUA{0klBE#<0dLyg>s2915a=pjEF_tgu~?7Zvp;EGqCn z12gqtgs|jw-tYIvTHl<^tYZww&s05azxc}_%Onm47t~0o*U;|UQ^SZH;?U+Yt5iMm zURjJkHzZ|*M7E9KYhSC-%%U=YC?jiAJmYIF)5r|R-sr$>rWrJ`HXlZM=rAC8t zD}vP|9mI7;W3_J(lIP5sGboA{=9YYmKf;eG$rf)Zd@jw8vI@?xuW_}L^MGIK#`^8O zs9~xemb8hZ|1_xMH!twtFaANH0uAmLso2+C83;*^NS)qdUhd9NphtvUAnuZf349wx z3pf1cKk7FN~dOZwAF_Y)#>V0<%EyR_?6_dPy5X{ zl-PgL7OsU=&n1x@(LFJ`KOuh4zO-z%5l3#Th2?U0Hi=4wTkG%5xvu}xo5s2xD?xAf z9O1Ky3W}ABjrt_})E%?RZC74bV1Mz!&fsgsZ(*D~oiOu_{!@R+v`l|^P5|Jf@7y}bO^Ej?6L4bd%KEj zDZbCY-q@kus3i{+rBin~jm`lls=~p7xH30b_hz>MSs=#NtFfc~fUg!wo_xSoTEe zz2#pE_+(8b3eIvFh!CK3-eCl4-nME)K4j0Trt87C`;nld1yM2+^SKQB_Lr5FrMVpKBia=qf1)%HVW!F^ zi|XPb|K-Pj(41)1|EW~0uxr1JcoqLYQ$b-)o5e7fy{Q{RZ;nAK*a=+$N~9WY>vYSR zAH^6)S)t%RXSh!hsulW-D<_4fKGvc+Z#hdv;GtAZH~ilUoeyjTg@mYBSWYDQ@7}#@ z*qo5sij;TIS_9SuY-6a{WS_Sc#edxf#%GnZjCH$LT-Z0s3SD z-_Vbp=h37bwM>M)~~9;c@BQmq_E1$!I9A_hrLgSZy4RD< zak^Y0GrEIgDG}{sKhp89EZcrBP$CV)%R$6#xRH_+6}2^A)YOuU2IV`vs5waIWLNbD zAGHi}>%MwL0Rp<>1k~Vrdpdba-7N9~`Czs0MROXz%Vuwq@dgZ}b#^@44%SK2JBikO zesq#_OPlP3T4F59j+YbeH0@6c5lVa?f9yPMR`$^4Sr7X170D3#L8je~_05%A2(pQs z!2quuA95(fomJ~tQ%>VCun&oAhX*3gHEWAsGmR8KlKWyMGGGujaE2nW(3~kK&YVm$ z&FT61bl{BAplvLljxYk9m5L{xE(JzsV)D^RNUwb8JyMI$?JCOg3Q6~TD6NJI4wcVws!fm>4`Aj z;i|EH5us-1`1Sv!ovTgG{g)e8+{ju)!T`?hdE?KvqQAA1i&&E0 z(oY?qyJm!GVL!zyHD5zx?H_WhlH9(J1*3Kke`tg$lwKp7Yo+oz5E&&$xuHS?gum_;A>=#pLST%`wjU3$kx5WY^c%dcBDDor$vbRp zY)sx3^mIEbUTTTSm)nvPS0aC|Ry@Hv;GefrOy5YFB+$ontC{}UnF-fDGR5NE9m7D# zMVk_F4Gj$@^QlO0(_v>lT+Yfp9Hg;(SK` zgC|Z(L6Dl2tSd=2L;2X@wVm_GU|GZ=x6;%ZRT`sDaPLzy8}-r*?e(Qd#j1%o+S?D{ z%rAq0^8F>k1nAAE*xCDi&or*lC$~i>Ej>+U!j{l~K}24p`arGhNxIY9i*LjxR+qwi zPtYJhL%9ZK^;~ZoY%zQ527yBkjloT`-E5b2A9t?m&21!RdDG!~LC+kl7 z_K=+N`H71C3;K4S&Si)g3{F4mFkcw6wuEj8Q&jvx| zAkO9rFk7(cbcEc_FiMKav=6PkBV3CMs_eToX|eXvkcY=vH zkn|EWBgp^MO1|g(M?NuP&^hcH^vA>tA!Jteh>Q0d5(WH!yIqcb<$RZZRSn4p0g_eH z3x$_y9x4yfWi5sFsxf3eJ$_>;+J9$>BgsRQdVelOG$;n^=M0BCpO%abcD-eKQ5w|{r$5GPHT4%$|KST)xag z2btrTR&9mEILxPwl9D>DzOl!Jl?FY85Q`w@1Ho^u49as{ArMVv9s)a7GdVc9M6f zUMsH|9S&_9Oo?A}eed7|K;M^L!k9eLBxJw|cn~~W7n#ZdtUR#{%ds!J|9-|M*OxQ4l{A$FOY(*qeL>{Y& zUkD=d4AGT;91<0k2`@y1!y`3lQ-c2rtu&;|`iwH%hJFQ}%zm)QvG#c8GwQySVxsaW zblLmShml6>M1GIRgOl)s8$xex3zp6t(`0YZr)|hk8h@S~LR~^5YmZMJ0LBrIm^dr# zc?v>*!uPX8Ui*+{P5uvKj>6r4LeFDC7f8@U654O|RC^DrDP6>f5_51Uw^<%e10Ztf z;;bNzh`i&Uo`-lQ>YPg1Je5JmlEUg3t66JG>H*)$^=x_;r+Sx%sC|0T z7`z-aUOJ94)B;5V=g-dadhCDOh__%Q+g=9C+Xlfk-xB$#uKihS%f;M=jMpz4Q$ z*>$3>RhOQCChc3R6tDA&0=RBNs~95x;MQ^Y-|O21*=$nN6|xPaGd&>KL{Lz04tzA7 zU|}^$R;v4o8X8h*i)8Ct{axP!<{^vcVeuCv+ddX=wmEO^Yikk=3mYE4?PlFsxs1gW zbWM=14TqFB6TCn3Z*h5ZX-8>QeU;7&D(gj6ZFY+7&n5eVn3h#IaSK5#el2~^vRr;R(=quX;)RCjX7Z`%`|9OL?;v)x2k~6UzKlXGI(md0%d~46D z;%W9>76}V`ff(@)KnJHm9sP(S3h;q^!F=e>z1W!qBoxcM+MXN45-TbU?aAN;#kzc1 zE|^MAT-IexwGwS0ZJRbs*n%|eKthp!>*V>1M(G|CUewvr7eI|6udFEkU@lnBv^d-~ zA5wcE$`lj!Zvl)erzfJh+6$Lea7u?9o>OBzH9(k_Do6-YB>_7Xh%7t&naK07x6L63 zpepiR*yQ>+z!RWA@nV`V&g1lZGcMf9Vjv|Y9a7SfSDIc}Fp^zWJ-j~Gv7LaTLNRJf z=$B6MdK&r`@l`G`2Hv7heu?S)hvEZG0yr&6YM$B3ofn#uv20a#wo$bXp#y{h!hT2x37GW8y1{ z%tf|VOMWAn{t?A%@3x7=SXVwB;41v9t1E-hEhq@jWT;pInU4$>8CipN9b`Ggv!=nX zKpv6{ZyWCE)M{`ZCl{6l7CNU|>Be95^R7<{@kydq5D7MAxGQbLQOO+1ZJ#d`K$3-w zytcQ&TDnztp0r;b-aB%${11biuj5yVLtk^nE^Y~)x1O0)b69A4x)10`l1w*j=X>A3 zy264qph^}7zhcrw!NQuG_onibfe4!K%Zc6JUNRc;zj&A&Wjy!2qLbb)d8Fjo+CrO~ z%$7}B8+|8!LDEMT^0O_xaFh0`F|mQO_K|jDm_RYlCNhdKg~l`P!Ku1)SOI_M7e%tE z8W_4)Ao&(zFP-*%gMIYYGDQYl9y~*Yf5;;iNC+%g8%;QZQort4?(En+huA=5=)+r@ zoVUCr^E?=X8$}__3!I7+;t^faG-;yw-rnAB6L-SWr9HPE>c~Q}PoJtYLeIMyhR|{10ggoI>IjUqV*=dMdivC>RvL443yAv7gufDWs<`}f~eg9@i7Ri=KM6%G|$4f4K)&G;RN8f&ON*R-nN zEogoe6MSM#-<;(=LFb+M^uQG&{+Ikiam5ER0u+dP_K(&;>~m5@OY)k|P=$>japSO= zUg(0b6a5zlREDyCCvN95H<}@RQkyHs82H-ajaDp#ibGU42q`Zgf_~$H=G#I?-%YVF ztKF>On(J7`o;&~D_D3H)El(gsYoR~gB#N0QzjwEXO!2`n;7Q^j03kfH_dww0*owBMgYK>yFgqlMY7M-p4#dy0Be z3@)h9qt}=F=P$b%j9lEPeS-nl>duY=+=*68MSbbuT}fB`B#Uu_LN~9_xpdR${M_yg zIrAAVQHko0gDam5d0%o3U2}^Tlcd`yf2zcWmJ0r$wCdg6YCfNeQqO;hl?zk5v{0ku$`zDK-%HY!I4THq=ZWsonk9iQ$l4*%RS|bHJ2}AA zC;`d793#MQHP*ybMLwbd2!T)46N~-kz{QP^C;a>=jYu%79Cxb6zFx}Z*pDyF^C;YC z=TfQedSnfCdv?du`Xw2x(_2^tko{z=7K{zXCud1C>2 zn+DO(K~&w;L^CUC9D|WuxOMbqmlPl6sks-#Q|0R|)FvtnTvoF$D>`nOJKWDkgCyC`L)Z(KjK%Y-u$V~|C4Fdb>SxmeI1G7q7>if86-qO?06 zc_M!SB(4*>hoBvMls`v85Rj*c0<3vvWKS0v{$;zky40!J;n%T~b%)vE5ETq_qqczwAwW-t-{Ue@43_ zIGsb>@3X&TK41PK+tfj{W3Dr{ZriY@p7nKg-v4ffFWaYCH|Mhlg-thqOG~GFAkZm#9%jm^Yv0G8f1JFN5)CS@~VN&bN8vH ze4@~0geQ3V{*nyP20qCQPYHtqDT$Bx2fLLAS$6|Gl*~xe_Hbo~BqX1Qq3gV?3Js&)$;SXXXBZpQpTHN2hl>jQ7mms23zk{|TEshV2WUdcZraMlmO=&@rdm$>)r z8!KpnQ6hWJt;t*aU_361vi5SFa%^6}QZ(ERw~h$QVATOV}}Je`Rgo!t?p)@%QEMPW_I z#8GW6T(-}}O}N<$iJ4x zAl(ApwtQNt<(`aCm)+mz1E;keb@v+oN1C%xT1*p2|1u9z5a9ffY2=-al4ey=;_{Yq zq@o0u|MV#3hAO06762J3ny(B|*HvbJ?e@=6O(+yGz(Z z)i+MhrkdPP9rFxvX*VfF7o}ewDJTzd-kl@9W@B|3^%C} zez@F6B_H%m&yo_$QRcUBrEoambOkhnu0MlDjLZ1znJ3xBH*|t?*d7nm{#^n^3<$sp z1O?x4ceS&mX}^87;eSSicRYe%ye~(r3Tf1Q&wHDXV0Dvn8a^2B4SY-;1EXjQz1vrS zLGB~|@E0M#F}((As+!POF!W%ZBUL(Ci&glngd2)Iu(J37DJ=EV3@*`CG+f)^7qQZg zUCRO-liVy%4(XTvJ_CXWInU|`S8CwZ`war9D_H*$-~xS4v_($3Db-=Ddd~h=C6Dj` z&kW99+-`WdV|Jv&Z*wtsA88ysZ|L&dYl2E@s|B(3oBe(QfNm+uLkW~!R^RXyQo+Ah zd=4FeNJ$n0wS)&@RKo|;WLoE>w&)YZF#EpX`5Fo}YgNNd^Ke>}TYl}`Q+@O?1+P;e z&hZW}nbT1&*l}3Wf+pv4!s`;Zoct~-FN}7DD9wCA@IMj}7+`9ea2_If0)*r z6?;0X^XkrLTzYTZQSuk&N&Y?ycAjUDBsdxn-~ zg3nn$dc0Tn_f}IhVAXEC$vTM+(s{$LgP<$!|;2 zxzXN~P!Y$$z4aE#F_RCROg(Ey8Qo}w>8U7Nr7EJow;Rq-%-q6423#pI8#1@B;oVj@ z?BKv_IdAzY;~_8>66%+-`Xy?%k3_PT?Ihbr2W4z+va)TgF`^O}Fhmgp4>6~e$UvL_8mg-i{FE%02?Nv0usX zFs*wp^~Fk=g&K0{KNV#%A`34%cOAgDYV!}-#m@m&>qT$YPfWZoE^VJH>iG~K*gY3t z(o^Ed)sry^si3O9lasFm1qAe=mnq(8md{pA17&h@Qc~DXuBElL9PFwduy+mNgdrY3 zpgtAM%yR8<<1TzIes&dC)F@r#(}ZPgSOs*+Q{-tqQnCvyvPE3t7vc&nMc`$D2e+AY zeUayu;}*r#cC!3W^xX0loG5lO4}Eq-R>%hjrpStJpX5apt{BV0oR*wiXsE;#ahR}J zPKzI{x6*UMR1qtdMtjNTBJhKtjDwYYQf zWS|0jmD8v2a;?6+5w)#yf6IXn&T(HZ`jYjh}Xf;u&}BH zbi6wBpYH9+@S~UpuZgcL_HBwW`q{UgDbL1qSU!JvoP=#=Dw+_oT#nZ!D=S-go#k>L zha4@SH_kbS`SkCMs&3_O;`x)GYqw14pk^x`aCf z`1l?WRi=f6aMX$8{SXxuoxH$9O-+6F{CRDd+lBEsvVt_}gO}F@c8)cBp}xnBx5g%o zUurzKa)2WpLd#`E6ZNJxAxj?0t&8XFPp9jO`2Yq3W(qOltZ1x`j!XOP+mAe}jhcv@7&mCYC4j4cA(s z+8Bz_ed`o6-%kb|_##<34puwHjK*0T2XsA>GZr}CahkQUe%ha-bI3wFL(=1{UHpky zp!*gjZN!JsO9O+0Ut4`@%;Cc%Oeb?c6dW4nTc{5o)IDjm;limmIyOyonv7j58FGj> z^(t7rdoMmN4(-vSNBzYnw?S>>1!{tOoSe7z_Vz5^#e|247aP>{S85oK)3a`@0P$K} z#@AsTTy~M&nkzSj*laGA|Jl@II);*SXGky{NPDn!m4Ewsi3l5Xpn7`b4)#EWXd%Uf z{EY~xyf;jr0H4l~e2{rG_sA0K~{@zn*;^ZUT2yhlTGX?uIy;%tX?&@R52{vESd4?Z3y zUJ~8@#6QcmPZp+0dIm@B)%mNRVmcGT!`{>WI&A1(0IFTJ%x2{mt?=n7BlZ}o>Gjq* zr&5j9()tRI-DKz-seO%C@LehGmeRW2jWE^=Ruk8V@8#VXsC#=p&#;H;5Gbxeb)E4m zx$_P5T@Xa+)sZNy-RMg=3$$wSX6M(hPyFwc`+JYOVdtP#9KK0)e6y&Q|I~u?yp{5$ zO^$PCdCw`Xx?Q}N9c6Ll?H9~0(X{I~iL6KXi#J_(0+O#{W1u6Qs@mrX# zzYAyxONLOrar?YgCdW8E-He7jh@@o3X>a4HhijkdTlISXf`(&@mW5^5T1o1q()gesp11S(Mjz zYJdMOH!^E;4>#euGZ_niv){2SsV{{DUd2`F7lW3k61bbed=!24?ine*Bde_E#ZWT5 zUI%?nHTix#u0Bb(Pu=qQM+?PAgWvud!^duHhsS<5SFY}MF}QMbWz)Lz7*AbD=)Frz zdl^{kZ$Pf(X???59IYX}cu`JI3(Mc%f6#Q!6PGl5Ns3kf=bd-&-qkj~Zw;Npq<)4; zcZAV@{Zdr<_3wlCxE}O<;`bzd3; zk~iPKTfk+xXE@h@LkoG25H?8~0>ZANuLrwp6CzSlq4%lm?d;+}iFK3Dh5L(Sv^S8H zMq&&fUVBG5*M^^Wi{`D}3$>!3y`kedc_19uL>D)NS$iN)_~4!)QLe>-dC3(K=J z7)DTvl0vd_bFY9YysKD7snqOlG5F+3)oxyB$m{AYe!Vm6T0>?RHwt;S{~y-c)Fj z7)@UkyLkf9V_0I6x}XdBdS9g{YUfLW{>e4Mo}n2A#i0W6MK$tmhTB8U#4D(a^C{cR5u=g58Hx zlapiJY0ni)UPs%EPeJkh6ct6|NDcRrLB?ZURM;CRyG8}A4Kz#%Hbhb*{A{r1_l6Z^kjpHmoV#@@g(WZo2hDT2CkaF0t}gbxTj;!UgPc4 z&|$$%bn8>CT|uukDZDhS2<7`F3G$;aP7;jQ)<9E$LW!G56@|vqTQBeAonjMxPkP7+Cc&9&~V@BpOai(l5P^NnCx!=9UH3Ygd#b!`GinKul$x|JEOpJU8>a{~DLk#RqRiHEQy zZI}#pXgMA`qq41EsuF|cBF?+^^UxSC9Xoen>k0^9irUuUNBVhp*{7>OuA37aM1gLG zYH`=&^qH+$MWHb%lsvVb%(W|d$CQjH9o_k+p`;E)8o}s#+!uJXbJs7#z zZ5D6394xUdAqpr0CS6_K_s}!HtaU>VW75L`Ly9-3r0NC-mHqD+YF(>2i?ew~5_Oi5 z?R(rL2l4DW_0NSXV<)-Gm@~~;0^NlaN-_4J{-ROJKMPxP@ODMmt=Ct;jb*}3CD~*A zn*E&Y<$ql3+9NHW+YMS7eEestMdU15^)aidZu_WOJ|skc!u?HFoZjlPRnXOP0_0LR zCi%AbsJ>Gy?@|Hl^z^g>RI=92{3!wgJHB-Dt_?-?R1J5+lnSi&*0`%R2k0mE`-Hj@ zcq@W`?U!4D!emTs4Ur}6XX9*FewkTXc9Go9c3i0{Q6fN{-cMBRqEC*~AuM)c0B&l^ zCGy{+CZkQl_yyT`D4Mo5hVp0*^+=Mh?u0DVKu@qc?J~f)$pZazK2kjzWz zFzjD3x8J{cBWfFQr~GxpvrsMTMkLlPOAHCPihgGxj@KZdDemM#@D~`6;lFR#+3+k( zZw4J&F|iXgQ6?Qve+jfJ>zE98(Dir7U*3Me<>eK@I~B{kvA&*dd~90UaQ-aP!>3b8 z=Ppsds^*e8pd7^1jmOBVTB%u6T_eq7Y7!QDQ@mVk=U{xBo0}Ug#$TRMw43?E^=VuZ zZZBYFUIMIUx1Ptdu&{WQ6B8Va|K!ON1_-pxl4b>mXCZj{RX#uKp43@FlN`ZWTsYg% z9{RXBp6SXl+y8bi$M(~~zdZKU;I8(Mf|JvgeW)*4YZ@baStA5+2?at;ncBz{pcNTcL6v1BZ- z-|%pYqw1B&ZJ)M-Z{)`b(&VgLHzE&wB#u8scofRj=3T7)T~{7?EKC@r`W}tg^^5~4 z#v$qF$ABlEH9tn%QUd6_y#dbH4j0yk!^Tqo*x*?#0a`hXdCfKbi=$y4QtXW7zx-DU zDt^CsQtdmxWjd3Ik8+AKioyAQfHb>vp!h3tZWb^XC-RlMefxGK_3gdeq!EI0Yw-Mg z`P8F1l_SiP0moOhUQI9Vyd)W8mB)-=`{iTE5PH(g&6oWaKO-q&?vb>9)&@yGd30$Y@_IC z;;;4fk03A;6~%EK*4*RrLA-l0C<*<>NjBd_rhW4}C%%h69~GQA6MJ?)=G?s-Xp+u= zeq>rLM)V^5h`EZD=-M8@#%_e`!(%u$_9>L_ga;0xzVl%!wXge+zVFswtBhqLIvEOa z%XF6=xVVX9^R-6L(@+pL(H%NE9LMd&I^b^dAS+2gSU52^S9*JX078|9hHgVmZvZL! z&lD-r-MzhL^#t87cW!y@6;vzQ{Kk)apmf?U)aA@eeZAz2n~NO?O%`pJ^=MZn$3(d% zl&h{so`3yK3%j{PlGQr5X{b3%RfVVqwrlVSFqZ2QnDBRx6eE1+KTL&cjWXuB4it?Y z!Cjq8ZeAMvn~IJuq$8b(h^X+WxdHPy;bkLtdL4IqM=JO9BE~o``seBp^!wZ8t|2Fh z3{ z(FNBl?b$!lWf&p{N3NT;U#^&TC;UV;Zu=d)7>$73SK*+(1Str+?Xs$$dw62veTZXz zq>H*K9?5#S%3-S$OIM&xxlws+mwD{vXv8_D85jOZHvU+OJ7anB&ynIdmO+B>Kgxp7 zu41veyhquRCz$o^FsAE`m_LX**`$?yU+uPo6E?x8GX5z>;hy6g5|ijvIPYQbchF@N zvPujJ4V~yr7V8K5kURzC#NGZ8%NOwZbBp5(GYg9+NZUy(jGze6PyX~*bswfhzc!Y+;)c$h^Gki7_k)XFzYoNGX`Z*uWBCKnU3h7M)5o+qgsyQ$}u^|i{|c@X64}0 zw3Xcq%%;z4kfkZY6QdK17pJMi6MOXSv$!@y*vxh&W5a~l4U8}Gpwu?AwGFKjpymH9 zI5up#AIQ*Na{=#LGY+wsLfn$Z8}v2)-G^?@!T@arF#B|*m>x|mA~l{>2Bn=HCw8xo`V>gF?*fv&@L+g@^S0QcjXXlN!d`HM9I!F&ac~F>Y8+Ua0N)bq)8@ zR+Ht#(7lE(Nudwj*Z+Cp48427NAatzO@i8AB8pw6LZCyu!AtpC7pEYJ;I8G>xI}AZ&zSo_j5&ljeYoqz&C!8Vu$H@>&G;lq zUPk6B+$6bReWFn=yZ~NhA!Fky8WD^~j+rVYuLq48A7zNs)Z-PPyp*F*ZMDrD5PyaF zzzNQbw$z<*SN)ORy1D0PKh&sQ>%oweQG@$l=h4)2!C zQqj!}!oLdkL>5__2_q}AgbN;@)YWLlO z1Q{vH{LgUGqB{_yIjsaghCk+o-_ABZ2h|5+(Nfo^qF-)&MdU<7g zbS5YX2)x^ltUh!)ir7OmmgSW-H8s8D2%DNxtrc$FchM|-hl%J*>7d8N#L@Ke>>`!z_ zNy+@s)N@qyXSGJrw+k^6mN6`-L)(oGmzxCJr}&#R(E5G@+7|GO2cFlZ;vwtF2`I%2 zkH})|#xI&+9l>Qkg;2eQRP(eU)+E?7e9jn611E0K@yJ?;7G8h!tvM`@P|yPJ%5S59 zQ1WvVF>hSVa`QgQBs5;lAZTzYtvTM$u=&u`+}sQ_bU*+u$QzTLmx*{BnBl(oJv&R` z_(3^!{Os7F7@Sn?7<6dzafq=UNDDQHw;Pi3qs};m4e5gThqcG=UJcZ}OSnrVA zQ=4{O>Wq0T+S-Ort(OGa7$5G1pK1ibVxPkQxhw;h1;-6~S{~2X9Wn*VFh4J|F3E3|1Lv3wV;b(tXt^@xeL4Y= zejfOSqabco2r$LhGeDULo*W-Nf?1L;5r^zSTx*qOTF>`E^r&*;{=PKbI@IGGPGY2~ zCeS~466{WR^=cYI+q-N$#lsG2#%DaR2yLAIUWQ9It)A8^SgFL?;)N#3>)HHF;A7ov( zH8bP4vttKe$a}EaXhFP^eqT1Rv55)Lkn(p-@Xd_RnAtIW%#8hL=Fxg3ZbFKGMceJ< zDYDc0eAao|j`#a&3?Jg3zIVm23EC8uAyb4(e4Y6ouhm%h=(?xa+tsPscmx1V{md~sO-m@q{2@(+&vmQvQ z1B=fD5h*Fox{F;YxR8GEQdgH;_)5{WXQO94_c5AM(B?99q}Ri+tkf{9M$->Et$i+C zyV|k!1wNhER5b&N+SOZrdm%?77GLob>018(Z#Ny*fgHaIw$6 z?D~ys2k^Z2Bz3vzdN4FvxC{T5dW1B+3s6Ip9?w!N-rD3wH!f%_skm0e)0$3t zep27)&0|eEJ3Bix-YTXuuD7OfZO>B=h`o@rv0+636L^wdu8n`|FEYA8r(T(AvTBNI zY{Z12;Z0`!E7Pd(w-BVPe-OCVAZA!3+Ue%w+86td@ThcIa3AT&yQ-!!c`t7AU<=Tp zMeTpbxp%L;BPb?@Qa7l_M)^zsqTHfX$MEVqtbg>usi;JXkRIz>(}F8^6%BX@^iWS< zC4G5`s{;IhKp(5W*QSYjWafqo^GMM&0|&=mi$N$%P8gK~3ksMZL9|FL%4wkU|55f9 zP+e_b7bqd1q)3URB7)LCT2cuE1PKZ0R0O0;TBJ(^2|-dY=`;lb-XvNjt-WFf z$)Nv6yMh~L6q;qZ6-Y^9%dX%WC!5pL2AUkgW%WtLmAjR2Xs=PIDtn4%SnL;bw7Pn9 zlImNMQ-LO;wb4r<{m%+&M?}Zkvvj5TrPb9b*>$RdAl&Lg1Xind6M@oP-SoSx3;A+i zCKV{-{XhVk0xgEjWL;_7CxS;FJ$q=*wVnfe>Evz8b~+P3M9%?GeCO<$HB?r9y$Xsv zG7lgAj5Na4_y3LPM@7gbb!26O%;3ff`+7x@Rb%3SnfZ8^`ZlU{qfyzkEErsK|U8!Xxgt2 z6c41(178vaS0*B7XXiKG-g@(0NzoAD?Dv(PKWAQE2AKQW-ImUp$)uWYEAb&*%7~Fh z%#W%VBKDY~4$ZLjd8JwcjHmeD=|{#yS6MikJtqNlZrblveY4k*<#7*Fp?-{1@2kO- zXPT_7*1csLj z+s4G|y|zk8&9ettlSC8Bgtqb0u2 ze{cF1pwyZ--@)*mmgkRDR8&+045#v-g)Gv05drX2oLDQqa76g0v+7_o%BY^)j80ZF zUSaa$CxpB8pCvEcka0@#(E1>mi9!=f&R<{nEg-zDMXIf1hkkE&<@@@V;tks=7C`7O zoYuZYO;=yu7-1@LJ-m-^x;C_kXROShTesO1#xV3wCU<)isPKj4ZFTr4P&5Ty=EAV~ z-S!>+8!rqP`2C&cP zXe@r967z#DJ-NlqEG{dH55`kIFuOn>t}MThkUl~~6o9>H=0hBOYZS5jf$hga4yzgF ztMZa|%mG-Jsxjn7(U-Te2?;RRD=-YrqcE@cD~Nw)6wk)=SgS-fH5*nS0MU}TjD}PDn#Ir< zZ(4g0XAg=OKPod7nGlq~_V@des@Dynu}QLVdUL%t{}cj<1YG z9PfJnd#{3?p;}pJ_NV2L(VYjs)I}Sh5z95eMW{uhQp5MkD@f&1J!<%lL*C=Q`)nI- ztT%6V%a!wJf*LdL^Nn2p>d878KmjN#`CR1Ek;9O zyvLC-I=7r(Y}(~!u|_Opi59-Sbuf{AZ>ZWMM7~RfyeK?(^&C1=zI(_m*+NLFCzsaa#8O4#k z>fBaaE5Y*CU)>1kPpDU}5NExdlv5O(+Qw2i-Wg@}XlWSc5QdDxJJC%GE#~~MNL#d@ zJP6Z&a_m>0Rva#-p9esp+vndSZ*XM9y$!h7;W!+QU`6!OT` zCxM1Vck)twAZZ(9{@q}bfln)Z8Jv;oL04R;vp(7cfjB>`loEO*mN|4et#kq>IP?p{LwLp= zm>!8|0Bm7lr=kP!e!ObIiv{c%#m`SzJLC8*XyaWrA9V1PKlmqN%F`$uFV|LW{d3m4 z;4W;Ejuy>oes!u7GZeqG%%IW}1LMdijYkT>X_5E1Vr7qoeig=42L`sjuWYVsj(UXt zSFb2*CU;#E-dd`sfr1-B(itSYz#N$FaL2si0;*Z|){Ed@IYsu?} z-|*C|t{3*bju{9w)2AR37Q1>r;q2_?o>kGI=Aw`?!J&dbo=U%O19yVeS z>ny~ZIphFCMDn3kn!1Q2**oZygtq;_HLFq!NqJybUFgr2ATip0@wai>$#zcp;f=0H z6RsSgQd~X?hvxgK$%*L9^8%Eqv}GKQy^DTE(NmY4qn^Mu^_8q`d1b}ujbD&@lyq&< zbDZTs`yW+~zN>)`{FlFetwWu6>cb4P;J)ubAX!Qo#YIBg4TIidj9EYY@=~IQL;rVb z(XG5yvTS)PlKQ^16WiJg=nWxlyWrzY)+_c;KLi*>+fAN)X=raUFfsYMyC*Fr1@tR% z=^bE@10N~rL`d-V?Uv^~4bE`JbJ1I@zp`g#Y8o4xkakRQe%D2xl* zPl)J5HB)ts&CP?^waRp%jG2eIJNY-GiM1RAkC3Pz5`%;zS}A5=V88?mm6rECJg}&P@@N%5hkosbdsGI* z3F_BV4oNu9P)Tgs@zgLh4c99zmT2(_2E}&`j0&lMKmnV_N-8Wq#+c zAH)1%aA2T`Hd*D{b;|pvw0?&SIYJi&25RC0y$bN{0W%P$lvl{%~DLXTBVBNI+6a*qnHZE@C{i)Pg5(`42h zs;z7tdQY!@bHR|6qqzhzcWvogV5nx|Zt?v2)_5z@>d(9#sy=)xO4GQ1#hVuABFk}bo+nwEz!A;*To&KotBdb)mA;w5K7N*sjjh2yt zr~g+VFH#*r1MxD+4cX_!Slr3dH!IFxcR`P4c#Ja%5e;x^R*RGxm*R99TM z%ZBLU=BrF3KqRO6NWs5eY|7o3ncR6of#gbt5_%KrDpr=Oz@p8?4cS<)nWl=J|V7ujkKWaK#I_H5Yf?%sqUg5^8aTNv{ zXxGvCt~KheQ!$Tz>YATVpsK(2;uJ}qlW!QEc~e*9%uu@x#lCr$*OLDN`%T%zil82F zFbN$JmIF47K-qoiXX8 z7eTJ0`-zdJ97lu#eO(WEXOMZI69WcR$!sY*2Dw4sT3cIdfk#=OKL}JB1!d*y&S)AM z8n#P)1fZRh1TwZBB&E)NPov-Q&>t|-FS_YET(nGNvRvgjRN_d_w>&5O3~e$2&9xKF zRSr#W4U;z9Zk*n1h90(KWP8Tl4Ee;Ty!8fJ6s&33ZD6J9 zYUv#9ij`6q=@#=QW8{>}rvIu}V#w9(h7XFgv5UmD#p7;5;K~{dGWmzn9x$ouDg%Rr zvu5R*v1x#lu=6gfw$tqUD|@j`nRj;Pj@r|1ho;Vwsaexwv4nhT@r(cXUeI`V^T(SG zT9jfr^!E%-j}L|nRG;0e5d1+_&(t^Vl0-A+|I77e#Hmmvn;FjNpdqf&qZKzEH=2yo z5!by@hg|nFw_~tGK>f7UWD7;8snIYEi1I6petHAk<ZT!f5|Vu)3h zURm=?O|?tAT;;i8ZQPVlQsecM&!C;cEZ(qD9Z7h-&&eL&-~Q;~LtGfI;t9n` zCg8fYwV0UL_%+b_i*Q3D6XdSI4&Y$*;TW>3|7t%~DLM0CMy7(?ZL6KXut62n4E0|> zY~H1}qf|NFzoPxo8DH#p;odf3NPRh;QFX`Qb$P>?;qUE9_xHXl?$MHNFm|1xx@^QQ zZg$41aGD0-S94yB+V@UpPNwwZ81|`3Q|2^%fuDdS6))LQ$MGDWYok=trH^YS_bchE z57(^LDBVSmVKhf!3G>XAO(u3vrwk3PHLiw)YJ$C+zaF`q@{QO)di+!y%Sz|UCX49N z}R^0V{vU5k3|@7@z^B~vF&Dj4cxb85GCDcPkTA#mE?FeBkR>2)d}CG zGJofcqVX!fsd|`6T+pb`7qDTGrxC#Y;9!5wf?{JlbK`z1TbM#3r0u5`4>=JGuERbO z^|UzIt92v89pBpB1)Fhh_ga%m^^9B_s$9*`cd)ZMovoG{W&bdE{DR$;Pr zu|G4ur2B`4HiG>`F3jcBtZ-D2-hrVmTwq!iM8!`C^$+S>h@%|U*ymuFmlUA1BAJEn zggV8Fw9>4QhrxQ+$>PwbM9+mhI#)CZX~M);z42=!kIrP?Jhn*jV%Nl;jy+Mb+!Jxf z$Z++Xs9uPtc0Xr^F^h&$={IHIjoW8^zMplCmP(W(pmSu(X(VpU;N^`lCW)-YLv{vS zieF4PkDu68HY&Tze_e4Vc=Bzfu=_DS|4cDVb7Pxw#)tPBH`zd%?}i_xXOG6)0i4*v zaa30!H>o53Uh+O0ffj6D*Ugy+gFmZHiX$kU2eDH3U);kut=hzptD9?Po|t^@4f*`E zxBs2zYf5Ksie#x$P=!H5JN9*jQFb!VWWh3e$c*5%3W<9H9mnO7{IzYUjjNT6pr(GZ# z)NKtu_A@I)^#q8-^~+c-0L%rkjxkK57wve(YttT`>oI={k$dpY{`EJcE6ti`i8 z4hc7L9!vP^C{t0e)NQI-i3JQQ_Hz2H-JEh2gbc;GnoBslUW$OhSh0UsI}gm;&B=Lr zgEpLl9Oa1*NcNH(U-B0kc|VG2ie(nRB5FZZRlL8y-z)I3LQuau1M8v=K7kTLy|Nb9 zlXq9{34gE0xpvx+zES5(XGIrq^pkTmCM!4^iy_j2QdGPsLcepZv8Nz4no91;$y&@% zyOKtLU_oViP;Ztn6B!nj=$+rO24U5|MIxTnHPhC#wUEnX_*~6UatL~LY#!}g@Lt3= zvu;a8*f$n07hee))jITb1_r>Dss)}i7x1H2yoQ6X$JaWHyB3>0`Nks0>(9|YB~o7^ zV7yH}r5T-&9h;g{iTC$t}YgaaDerhoabvFI=8 z?#$)xrv{Z!5bQDk%06@sYD(X4V~ZrmWz$oi45p9IpAs2dZ&lg2ZL;tEdZX&YpR=vG zGe|O3tizV{!8zZLLo4p*h~redl*=Hw)V*4ZXF+?Xi9W;@w|$7-PF=!k`hAQ(zn-dD zQL#*2HU9C<>vz_aSch)}rgie~autX21Le-JAT2M#rb%Om!?22%982TWH{b{LV2U$| z&92Yr{Jq@N=>)6X0b=>HS*byQie#r~_*5f3kod$a4<)Pwk+u?#Jy2m)$ zzx_1Wzx|}%VhC){bqq@DgXxVw`Vo!(N+^bKdSfGwjOs%FPORGo<2{!_lZ`wFOo~%% z>P4buTr}v1D!O%xT)3r!!#l+zBR{&ucdV2+?v&Vnhkk~@Ow2gTwUm!Fg{y6o>ohfJ zU$OQH?hZGp_2P$7W>KNdt-W9VcZ>J$n(te)?MCsqWs|RoA#Zuk%sgGNj^p$|Bu(#y zb0~2&59wOBPv_*6h);I#4$1hIWPA`_$&}sCbsFiJ9?tM=L4!Y=^c;cwd|BF}Vyz!h z&+jFJW6@D3^#wpLUy9xVllaJ|HMsHrzOi&Llq@-0Vwz7F&NhYaqw}S6V zK#_}i#qC-P<)uT#0a8+O@+ruVkuEm^YabuG3}ULOt0zHOA1!Ee18LuW407W3`w_hn zb=RwGN;!NVr8FFMcRI#ss&w=k-el^ac);oo`=YjPl$hmnG)sjj-(?Yi_R1~kC%6(@ zvukumZIXIWHD_0A_|%rj6A04XZ)U7Sp{{jSa4g-xHU3pis=^dN<`VifVe&EMv`&j3 zWlcl9#qW(p$=vcHohrY%^8JzaUno`8e%QGQYA7@Ms-BT~^~H97M>@XKR9W-7HpSd* zb(?1@NepyKypLAYR5y$bjvKDIh0~AK5YPpj+nD^EHQqAD`iL!bs=Xmv&p5(yDv*Ea zFn*t3MoCnEpk}?nK$xK4GL(-D`7QD3ZUOQ8mly&lrnE^vJSFt@nX%gNBxMy*w)-0Q z-kZ{RXR~1@YllKnCb=lda;+0&21?Sm)UkIw%&RqYM@y~8zb|&0)GpEMP*G(S%@s_J zk<`;p>z6PekIesg3L~K+v%wpm@~a_KsL<6nDJ{s!dD|T;;ZmCAHWcX?N_>#cHl|!P z!c*Aau$bU*ZiAL#_Z5#@J85>*qX#Kl(GHq>=H`W~R3A$CnsjMKq!+R_Jjz$Cznt+* zEOu%oz3QloW=e3$d(iG?_HZv{vbB&)p_f=DBX)C+Ch6vEZc(R@&)%i!{zq;vGZHrD zBrt2+n7f&bRru?5TzI4QqJa!`E5Z9^!`#P1oJB@w zXZ*l!y1&(Uw6R!c;VKmmLp_-o-fL=U6&?f9&K-JVs`87p*Or<>76-ylD#2Y8oUb2} z*P$ngVvbb5-5L@ZxQFP7{>2^#HZXf9v_iw{Kaoyk{)&PwL&+_y{QEl26#_TU?-6Ze zSL&u6ZK{mxz`T2nBbSOWU;4I&w@z|yJyFB9eMi&#v%aB0A82C(IeM7Tf@Dp31(SXd5+?xM zUWtkHP}cs)E5ZOB9pU!5-8X1J)KZKBQK7Ym;p92TcJynfIQV{;ua>c2pJ5B^MPUMd zFmL<^Xuo1K+H5O)V#CuH_vO{kUv;3`nTzBM3d$&b;GlkJG;?;!hZ>^}d>x#SB%Ezs zO=E+@KM?G4{$*#T_sN@Qx|3tlyN_iYPLG|yjuIOh0#Y#79hi*&OdW`MpODX%_2++D z^XPHF5EAjMt@QvhnKe-}u{8B zaTy2#{d0ZaU+#Nx<~zI>>HX>N$_N5sfAG(LBAEN|q6|mfQQDW?@G&9k&r5OIwESa3arnIiUtg!ksGxA-Dz zqD`VEUUx#AW8zrHGs7I>>o9%In!Nbu)S9<0iLDmmHPaq6{MJA#dcnR=;0@HH7lV_^ z!ji?>j@832q0P z6&*(74UFMsPZDmAH)m+x(#~90O7PL|o}zExM%yk&D>%bnkIQoYo^gnWg&+#Fe-MS+ zD}qtMbL929-KT~Ts!fR<6P6TEc2W0~JGc00gGjLH(2?DmLNxSe;FGCvR2*84GdM)z$yV%Gfk~z}Y4_Bg6VZHc(TMTMJ|(4y*X4rlx}uw&6w@IQsGE zyAIg9(rDfCXbaU#w}Z6LW3dp`%n{OWYSis7V*@xnRQ092)cFtK&>NSFU_QK>vMhDw zrX15L!MT2&8Nv%}ZNr2Pwz)MDwHOn#hN<j8^2Qtlibr8zm;YXs&NwENMyxw(;5k~U%j zn~Z4=`cs}M1|9++s%bDr8bo^SclqLQr9V2y>!$Qr|p zrx);%P~#?~veR2^1$2fEHoeO|-CCv^#+s62oBeMOR(n2Pl?crzW=p4Y-DwX)=VUu2;jeO z?TmrKduT*DPd@Lb`q_IT3%$xRr%1f{;H2r~K(OTWkb;y(?bf{o^+Ty2dWS5QCtD8( z2+Qz{Bl+@tpO!^OUp@v{VW7%Y=v_F=Gb&t!O2hw>8pJoaU?RV721~Zasr3{lKR~ z_`ZNlcmUrw=+Y<}*{t}Xh4}U~A|yOzuPJ?I$Kgl?>+5s#Sfld{5ySM#^T;aeFXrn+ zBm6jvhMeUUFqG*qX*I!L~FMDxtL+E@!4bRm)X=w z&qhs1<4*T=1V*_AX%OyF66*2#8_whBaYyKPve88B(J`ky{W#FGOUtOD_gG;2 z{i=gzy8=tz*22}ylZYE}kqp~r zuiC|be{!$gh%DA6`>7%QHMO<9{a1itvf7ORW9XnlB+Sp303;EKga9l)G9}S+P}lI> z;s7ovhVxyId?L)Rt)81LJ#M_@asZ?PiMEWQOdzrGCH=mIhmE@+qE-F!^3UGG06yiu zt_16tjyq*k4-5Y)(MCZ%x54n#&jpaLq8@&FY;xmm=>OTwolKBj^vHjzD zhQPa946jmdmrSmBI$gG2q%m~mp^|c@XrUFsTjLavbt2B{%(%yiga-pofB@{Dm&z3%dN?sP4j1r_bJ)r@dPoRPhZxbPV|>!M zFfY?d)VcEqI9*F>e6y{N#UEx8p@xGpj z=pKu3|L=3Z>rd(8++7J7pEg?O_n^sfAv%%2%`TwZ>c3 z)>UTbvv%LZy^Y>nrU_@di@RK1211;Ydms8-S-N(mTgDkReN|4<5{B)Gm&Uj7+SpWwU)Zdx?bt&I7lg@4@FCs!h{=y%YIdD|ZP$ zb^VZ0J5{0-nSMS~HyoM2*aBWoJTx|6u(yz3(tu7UtBW`q0Gr}Nh@f%`52Ca6#SnJu8}(620PKIA^XQ!>hDRPj#oS zL5nd$B_`ge;VK(F&U1Y}O6`HlEes12ny_VmE;@E#y6?3)0`))}O0p=s&7 zzGP+n7bO&yTT)GABnJ^Ft#(vXT5lzWE>gFr>@;Pc`tUFuMoal(RU=tF5(v|{*d={wSdhF|~YN)g=r zAM~2w-|>w!-ReWjc{9uUnKbH-_!8EANiyH>WSY03asrEVVjdMgKqUj)jrYojdmDV_ z!m+Y(F6tCQLglyen=>*p4CPzvk2bEnH&!99$H&U6n9j>bH}m>y4Sb$?`F8(@FJ#kh zAp6j>sQD~xKr?fc+r3zY6^||8=TJk&51b!GYiv0Im#;A_f}k*^N*IP3KN+m6&d8A6 z?VXvNv`W4Idtu>bN+!-)1W|?;t0Xz5BY>{jJZ#TqwS(La#mpH4%>%^E%`!O{%b2bF3!bB zT3uSY5R*E{Z_vPpF$(TX`C^{w~TTH_C!!{q3V3z793d0Q{qSOWE@tQSM_FUXw_eJ=P}ZH(2$ z9AN0Xh}~?p)gwCcaOC9#vZ#jN9LE5#(2AG~!oj0S{ z#+n|QbpM81Ksd)&d$v=zHnB{O0Lo5!mIv=Ys_MPeC=^fADl3XfrA};|EBKr2^m!j+Jp1bty7PJOPM$Y+qiF=c&K=_uYxk&QnRELO?Y$m3cP*!kBa^G{FRZ2; ze>Gkf%Mxw)LSzh6g02*NIpSLTQpId6?3KZ)VXy7&ZDXX+X8Af*qSiQcLooBNsP`X4 z(}tP<3^Ym|(@e4Y#06`WPXmo2-NKMyXZA`iJS;YF&)}qPA<_rJh&mnI4w>RN*c}`k zC@cLHqY(K82x=y<)3s$y#XZ@18R+82%!3+w`Ls`YR!t>6X{m+5Ech``MH z(UPAhZXA8bnC_-fHcdj$->S?1qAmOfbB(87nEZX@bmMGn*H$mndDg5llj0}kRiFu4 zB4jZ7WHNE0rUVwY@2{zYuMp@D&o5udu7ncUm#NaIZu zSG5V+m|iq6Y?IyNk|C`NlFuPEk61br$!Zv?{%dXN_Cyk*pZ_M({gYV(ygMN6{s*D$ z`3tAECXV_vI?={;nS@rqc)QOELvBVYS1Vy%lPBj%Nl7(h9^J@#_wEvGaPTpG&lsg| z9P|BAXG-+1>yf$v`s$NTlsr=oViRZQbt#glVzbL}w^d^%b620m9Dl73rcy!%6)w&7nvVE0yf`=eF^tJhVJ z0rgFdK{5NQE1qs`N14CHA7d^HWoVEpyyeY1+Ibt-cHw5Or9&b26{TUvgP#43cH_@S z?iS2Jg3Z@ro){_YNKZKj{a=m>W9S^w+MUU4i<-~#c|m9gv|;P9N4TuV*9<=`ykOf) zi54*iE>42Jy}H)l`_RPLe&;0br6Vn{G4%&sC^_{>5;-EU=Gk+TT)2Oh2=itA<1^!j z;-UM>y~je;vZ7>E57w-=HbkWVVX*PP+$?Et8>7c5YRvn6^HQV81Zl1opCy>_y0RU( z!G%@<0b^N7g-%CjoQ1nZQ!Bd?&AN6ox==zao%P@zCFfL^dIziXee4f(Ria3G@z>}1 zuVD3@-5q2tcFGEx9(Z`CimAFPFO0o|COLT0U#rV7xLI5`ez$@MXjHT6EUHg@zzg@+ zQ%IuITOdmSXRh2yt@ zUfx9f9!v(hQe`NDXoSG7qD>q$PY6v1x)067j~@Rlu1Pwr*|;fw_<}M*I1NLqA8oW4 zL%X;M!~sZMhY*93m%dmw@;9Asq7N9WuIkv=m|k+BJH7Hf?4#kEtYWYTgrxo)!uto` zp%5VXU*+Q?H!?D^eiP5^7unHcUV4DfkZj{)PAmXUzo#aZtt#m z<`6-o1^#G;Fd~zMhZh^!cDNC9M)Z51pFPewHK0V_Y^Pa&KWD&vnTLpo0Pn5Bi!Xp* z!WgwlXu^UATx=1=AQ)BVnDpELQ1dbchjxbDmK!EqI|wld;H^e*dg>~+&^d>V4ZQWc zAb*|uAOp`rKg-0`H=dLPYrLoN@Wh^XH<#J`6PCW)z-^Th~kfVcPmC5%CqfUv?5X&?kYU=u9}6wI z2ae~yd_5Ub%&p*LMlwKiKF;-WlMuVm^|N~(qD&ryVrcAMi1rqqED8i7$Wqu?U>mY- zd-o+1g%Ovtl*QI4{JV=vKIEmB-inry{L*6u1wDmpk-yg$u(3m1|~?s_OFZ}Fmi&a8Umbw2&FqxEkKDKCf4Jq_LG0;`4L!CI=9WF z7w=BY91u=(JF@KTcab|5b{!^A%64hMA`E%7{s6aHUatwZ!j5Va6yn2gG1RQFbFb57 zKYOgd7EE#DT#|z&w@?_=5K&PKjEq>|r}Z6-4RQdfIEFC*YWcp;9f)Q&vYQalz&cz=;z8LrMqQ&MsnN0gr78va@>8W@9ci9qcNQnjPGg_K?NnKki>LIbKEK zjmh;n>giv+eJIp47MWC1avAKI1zZn#kWp;y)QB_#e7t=D)7sR>qwa4R4%h7}B0TRD`I^ZK-+%H(w z8bahV@N@?w3YzXGqzMBathv-%%W{v5glgWU|kC(ah&l!VjnF+WQ=Qr17AOfrb zkdSw0K3-IuXnH=%i*aB5y4SlU=ER;$iT&sRmLdKJzzU4@NA99mGucNMOU1+j1CCqt zyR+GC)=T>n!ptX=+gV9uP?d2Rh8A0I6B0s8N(5F49$LMa1zKD#G=9v63NCX2eKAGr*aBsS#Lg?imLI{uJMl%|BVfv z5yqIKd<$7S9&`AAuPDs}S<%yoR2}J#*S?QS%sqq9oa{ySdyfP&tYXp%kIo61a<@^N zu@=}Xy^oCz1m^5CumA=soVXED&%XlH?CflF_Y?o{o1>$n#Gvl^J48KAl8rMI1kYwk zY3U!0jW~UMeVg)c(<(0dx#A4;5c-wL?}<$&F2DV6**Wa1zZ!hMNa-|!m*v;Qn@cNE z7@FJEZ~8^F>u(NcJ=KT{IjEZ#k2c9nV}QZ-v5fX$LuEEnLL5lJ(EtNR3!s)$QFn0> z#>K^*ng0e}O%ou-T3@%^UKyr9+%+c%Q5efFZ79Za z-7h=+J|tLTvfZeGP*&K;A3mT1=Cn2l)fMSV=(rtVAw#2g8HB(i2@HS`>ghdIIGR{+ z%fL^b2+?e)Emz}ZDHa$ANF|86UL*%(jEKlQjgnXV^a0|J7qwbdvQ3oO`+SW*77Z9F zCl8`m#?F~}{o~>@LayM2h}5`y;8&4iIp=`FP%bI_&Y_n*fs=2A!f0Ea^%gojXzB~T zDwifJBO_sJ%L8NTeG#JST!MR_T;52fGe0=@y^V_tg7V>j>04fIu4i^O2hw$15#)2) zG(-|WkTFO@p3ns%5A#sei#HiGHD8B4T|av4AE1~8ZLxodVS#MHTnCLc?(dMsivvqj zNLRIWB+C;k$4sTTwfFIs|79A8_V)JcdwczI>7<`kfG9x*e)m5>4+xhd(TXidsmLv4 zBa3ha_|OG}FW)~^z-ElPD&CmjG1iE7=KJ>*@PP`hfG8KVPxb|4iIGP>gwo4M6_Z}RH$%pa3DRv zh8A)?SRQh^9@-*p2^de@ub5xj-+SfGE~|9X79bjmn8jR}zm@YL>|V zK9au!;7GoB-*Tz*+R>KtW1u5YqC?Z|>dPr#$D2iBn}Ivrr9hNIykss@|4uG^0>en^ zKvG=j+Vm%RTb!Di`P6=CX{=HRG4=&F1pT1`lRw@atiW?e9AkNIK2ad7b3L~lj3zfu z0F5u2cmJ?+LNCt!J7H+2f!wKgk(5T$2Z^0v-h^&eCf^XIkT5ac$nq>)a9Di`sDty} zeGjC2FCK;ZS2iJ*P5F&|wV8&zugcwmH^w+7i#*CXd(Nz63662n3N$ab(mu}d`C7p-q?cw}Z~9)Iq(Q$kYP^3&^W^NpUbxgoi1!#(eEY2KOKPZLdG2P}S( zr&kfVZZaQ&AuIG@*2DpXMB5FIuZ>)1S|9&tg2LD@RWrrPKd-}^U)<8#S{QunkR2N; zQ+|pu9`ua!*WLr!EmsV$5wSbYzJ->>w~4qTYWr9TLdcjVl?w=WDRpY5rtl8ie-ja5 z?!LhoHrw`&^Q!bD>G8}D#fUmh)wvyAdVyIW{Rj_lW?Z&E+BN=)_EkB%c+Lr-z|&7np23gf&d*Sz0e52~ z0U}iw2aPc5(!en7fH4uG-rVk5tG3(KgU*dd9hP{*EBdp7)o~8h)7tyt`SMY<&b9}5 zIQru_w}Rvt3j1eymL*(sS?ABt{&>0~FPM!CS??5(Y6NI{P%CjkXnd7vwX;cU8z)8_ ztYUGuQpBUg?`7)#Ue#q}=uezSz0G?Zo2pzy;BwS^SU3z zmQ8 zpg^pB8UM{t%4aVAPlESBBtM`m#E9X)yoU7{-bY=vm=$|$s~dwqF9^LKXsAs0v|9}j zYpIhij(1RohTunJVmFCk7_R-)650`~&#Ovj&RzHN9@dR6C&Z4!;c?^0NLF*IN3Zv3 z%rkd16IN;=jqLC!War>JWMS(kq&?NjF3_Xt5`LzgdW>J^3cA{GbcWSw^fM=qMnlTS zp`E7x5=j1Wb^9M7*8Z&RO|DE`l}NQxCuf)kQ_|@~KqeZtuX|!dU8PwnY-RX~YL#Z0 zNu!XN;~Q&;e}Z~hLh3?Z{!2>6JIn>E7c>%xElq`pbMW4a@QZ^=@Y;ux0*qLNjEcvE zd|}>R3dE0&@xPdvMlokNwSMQ1iMC^kEci_mh0ZWmybx#v$@0YEsVb=G z1@}%_ewxps&hM}!qz+B&dTD8E8^Q2_KDXI*=f~5XK~p0q>m2_!^K|m=Ami*O4U0E5 zKP0nih^5Ptw${1p*qw%=p>Lrbg`g}?Hk<%-{8wM&KbHNUD1MIKt%m_`IU)^cwHX<1 z|BkebKD-9NV`E3Hf*P@b|3J{S!3!fl8>y7aoED$anB+GZhOvr04Edtr7DJ^NIBe>T zIw8Y`u6+Uc?SB9M-C7y%@jGl%Rr0={Q=va600iH}qU$k}MIBygqa- zEb;cnx+@_8@^Ki%O>Wx!qLQxO6TLyI;7WeKt*VI+wb`{?%CqEByj89I|4Yj3L$wFZCYueh@S3+uCpsMCfp z-|7Lu_dXnO9Vi*qi%!v8N)M0WZJ|E8mU$gxPSN#ZR1dk-tX!UwJZY0;lB_1UYNLiW zc@N!KTQi?Za(}2ytF~RiCgq@N&;4Ab_S-ode?;^D$D*{P6yl2p-8Z^_p^}{y^7FiF zKb7mQSfsSmv)msgEPvONS4`Vib=brD5$A;rn>O{eRt$DW%5`BzrR=;)Z>fJ6F5FHJsfk zhv1=)%^7}4aFCB(2ZMO8>(I}N+jz$MwGP!so_=TOu5!sG{Zes7j+oUnOMZpUo1nL? zGd@Lf58!V==e$NmgR^|KZ3V5|vc~S*gbhX4t`q&SL1GUoan4Ow9E0=I`FRSkxWTa9 zN31|Et_w^El2|UL>m-}JBmhxinfXtxrC3cwIo^mpt=>FmfBw*JR_Ck=%Idl$jkr9BB z55>(roA46`7B)_N_u{f!R7Pr1NKNv5UtUqTm~W0#f{z6bcSFK>Z(I|P+{-V|m*ULR zrE_H19b;M2u%G=Y??8cy&=Wrtz2+AG*rwg9<70dd`x%=3_Xjz)t-mS*erf7y z8~*n6a$;DwSxne?wcPK>AcXtU%l>NTiF%FAZQT_wU@+^!K%)NjBgU;|JuIgQbo|Kbs?r8OZ)F$bv7qrvLTe#)yM4%UJN)WftyL-^@A7ViDUq5Rolv2D)4Z}VH0KS@@X%1(;b?Y6& zc3s*j;(h=oM8r^D*w|F@`~fbB3+<1M?@RK_-NApmO=D?KYnh0Ha2fb6l&NRGckDr7Cqt^l9)>-Fdaj@TfM}7l=UzRUcTSS z@?Ffm+|MAl^y%=ZFh#g)xSpL;c@ZhMP3 zU?}p}!Seee3O{z(LWd}ZnFDr6Fb+aS%fM{e=C1yhV~ajwXYl%r(Dqj_WgO1-2qpln zj_p!TNxaro&eT}3rz~e4t5Pk^Bq7DsIOPgu)wTDP^+iSZS9mn>It+qFt!#^`@Kv{5 z2Gg0K&6gUbZ{VIw8lKl5vrC5V;a9HR+*}i)VF7`!dG;;*FUu>c?Qn+~d>)&m^FD<^f+(&g6M?MnQ4fgKDOn{oRXZCfb2ykmiO z>F{I@Jw#+S@dIw3d(a47g0z!%SrD2%1KsN$nzh=IW*Ji${AT)Dhh4J*-|90_O;5$5 zwh`j=IGF#LFFgaKQR~~=5h8zmP|5xf_@ZU@j0=cTX6PYhm&q0dgg*LZiJBumP#M_w zW(qS9!bF$XuQNeUuqxGDPa)^vcCLJq=*0?1laY}!fWOuv^3bQsb{Eoo<@HXH`H^BgCXH ziw>!1!A05tV2BdSJ9`dSZf4}%4KPB(4vF)yO7ZzVKCZL-SBNA2iXi#CnqvyC-&yw( z@#o*To-1%w;Jb3ez5S>O92jol-ilHi3U(u{H_)m3uEeEQ;mCzZ$+x##WH@xzPUWGA6`aZy;oozVfD{VH;Leb=l@WS%{w7GhM-8p@>^@-5noS>U4jto`MoF*0}B9_-)7Wz`9PZP@wS`VjFKqmD@|%`QFG{$im~ zhBn%V9?!kKg57j>Vm`zE>H=B)!F{Xq-zPo34qQ4!>xghpNfbY|Nd=Y`J~fKkw?#WA7U@6MpAn%bF50O z#@Cljj8P3stbX_8%YQrHs~ys1V+(F%z;uR6|Bb_f7Y8Bfl%Y%fGM4e~_K%Mzu11I{ zGd%N$2C(iI6i0-XzwZooSas&e#`s$OWAU4?B69J@L&;y_wxZ0yF7{;$=o<#PMT z!JM3K`gVCM`Cs3#JGd4;;uEs1XFwJ+l4vxKgFktN3A3=Zqod*b3RUj>5te}3geU%y zI7DB^pd~3Ii52T*{NM$bmc~JcN@Ex5wmmsw8QwdTZFLmAkM1 zklLUf2sYK6{MmFXWSyx zDYj;$=hO4Yv62C_rnuYo#hICzcR)i58}Z?mJ+Xv))Q{at#+bdbs>)ndtAhlV!c%;%nH#n<(xK+n3Z24QcNU6duTuf$F%6`m)zmuV)(%4OOkn{Mw>^#M%eQ zG&MEZMMO@66j410VCFk6s^HMXI$Q0JIY&eYMGy7o?oPXQQLHM}7!+Fy=W_aKsF8Uf zl25|RPfyG|FzGq86%<{TqS(B82aTTeHJ817t=xHPRq@LYMAb7YyV~zyJug{s{CT6u z==Ju(afQH)(+I245?cZN0UdxZnOal%vhZ11d_>1WE|w=U{NPEp)Lx$foF2NZI_^r@ zX`-O;y7*kp*S5B{^6gTI$L)j^_TFs#+WErx%x&c?X%3g4gl6-0ehoY?KG#$yy-33x z-1qfi+Zbva zcIiSb(F%jweQ8?f^fVK9FQ%;i?$!^Myw9-NVYiCTM{~lXk)qZO#{m24Q7TPOBA~VI zEqcD5a-T9CtEmDzh|~V}#QVQGrGRtr7X~IdeL=0mnFjGe(>@=OrGB^HXh=_=cF;OI zr|R>iAmgMn^eiltIC}U-A+;+N+8_mD=j<%+)9}QfUYg|lcH4}W+9dNFo6;>g!$Dm$ zVyI{c3TalZbVU0SVHIYe6GT1Ke4absLA6l9K*~#(FKlRX_k^K#Tr z338XIbWLC8XJ#Dc^Gg9@+jk(P_O)(#7z!H|MxXy9*lIzX!=_K)E`QqpTzjT5QCsjFDdZ*Xa;!;cMroS zP;#@iJvW`%MR@upf;RhIsz8z!Ul#0RR^@Tpd!n|l7KVf*-*1Y?Eb(vY9A62S%LhtHsiP8Ucjod(lp};!Y2QQCH&jq$=lS3=9xYRKs$H7D~;a5MhLx<(pq=IuBdl3 z^=@}R!{B?7@9*A!cp10KOS`a`l*CD`J%5-d!UMm+lv7$yQ$|{8%NX9P$x?%HYrMeY;^|hAZ(ZW!3>n1_eledH`gwe_CaYFH)`Vw-L7X}uUdpew*`FBkE&gc}zex_@~A#_3|+`G&{*$qri!iro(Qo5Y7 zyd*tt&{^akS%=nHu?m{}z1!v1csSB-K-*NHRT38$clM=FFF#lMLx++3gOTcI`9Z?j zj3GjMJsn!Os=-TguZ#PpTJ+H?%~>M{#&~k7B&FjUgmldO5T_& zm+-M8(+VavSQ(c!y>4d{iyk%yWiB3NQa6er58Ra`rMXY8d9IqzEqwhaGv!^g)}8*% z{>iqP-mDLXQHd24AmDAq_vUp%(uN+eX9Id?PTS{N)HQi3pV5jgy3}5< zXBgk+roJM$aMGSp)Hl6=h`EEwmzrksjxGygn75;ooBW9u6?fiHkmbzQpk_CsCX@2R z8M;7-SbtM~x@o~?!o5h7$A%Qw5j4;Ci*CDkE6vT%8F8C%(A*Ada2~qGDP1KumX^hfDOgSCs##+nAz^z_F-XOTx3{2L_sSmoKD+7EuJo-_ z+0RYxt)>@qaTD-_eroNNEiH*JzS!Sc{Z&ahSL`25($Z6Iy!`yt>PN!-0V(DS>{2`= zqEJ-0FjK?iWVO7B@=+lI5q=>j;XUGd(M{dD*OY?YxexU$J39L0udk`EszT|6PCy@T z5cX-e8M0ts`EDL^@^TG#Yi7<8-fwu4#Qn@__A_5AEjsHe$KPV+{u?OkSEwxUd?+ZS zdq8F~NNzeqBPXH(ApTlhL5%!D$*TeBoRNrDQ7ai9vI6rvT2SY5Z-MN)sRHo2aePh% z5|u7|VX<~K37cd5LaPx95&HrQ1NV&A&hlOjL6!tJ4Fqj(`&eJ)`v;BbeJ44jeL^Gz z29FUX&%lS~Axd5ahhG446P8tHE2g5-fwEx{M`QrNsigo0LvF{%xjb^hnVUC+x#4h6!x>k6r<=v97-aVXamBp?{{H$IYc%v{%;pzW)5h6^$*~XGDfRSH~Rs^PV z{)r?{$9-lI$>O3S*ahDW=gTI$TWyp}8!W4@!or>wIDL1=uK$VPMDCg#IMB!24tX8o$6VS&>uEXA<0Kt>m*$EO!=E}O;^d#Hcl-)@nj zx6!yWwL^q(xOLp=qY{kQ%~)`6@NdDpe{;yP;1gSTJ7PBFs;y%aiQ~~ky8blG`2|E2 zHXA{JX{3GVtp@yz6Tp2xIpP9beZr*-;EZR9j?i^2g!PUNmNJBU0}j$lOcHQ@)Qjj9 zPR5U;o1K3Lmi+gvvse9T$0(UelaPiOQcO%Nw}g|PfdRzOPFzEw#A7vmhP4f3%PQlq zlZa0rlHM4dWdd?Ey*w`q0NDiS1+5WTj6C4$9!KiS=jOCkMQxVxZ+<+^H}01=5_DX? zK%mJ`jr{!YA^E>qyfF=LxsLZ485%N-lEhP=hkJs-IH#Sk65m`nPLpkXNyUy9R|DKP z$1XmmXSueLGB++AVG(<1EAJ*>`9|9m>XmZ$(jJ!x!X=DlctMfMbPDg$hcGo3P=M9k z=!4niQ!J+SK&bsk2g@tFOy2?Z?PzO6YB~Jg9ur`7QyZ_$%`Nx&d>dmzAgaCK>gFK$ zLP{C!lB3ibZSw#c0bA5&6*cbY7+nWkn@3j=!k8(!_Qc)&+Qyppz>x0{426i54CBYl zud#3c23h;x&R`4Mj}ZvK4UZG57~KVO$L$vPHmcD%qNyA~uV06SE?*8ZWMyTw7%S!j zr~)4xW&{+kM+utX12jvv!ORw@7wXM?AOOH4mZpsfLPf^TRh?6`E-A*BI+$2L78`48 zzW`;7VI>FVW=E!+AR>k@hmHjNiqXi5y41yI;*-bxa3Szcp7WjL?O1 z=+{aSV0uJa6pGL0CJ8dg-SLuJ@9sW_*{vH@?akqAgI(C^Fz_%L+n^VpkYand_VVFY+ppu>utN510p$n-N*(c*U%u_B^da70Xp-zDGI>xWr>EZ$mtiwB zh*RSP_@q|t_6yLLo$=Qh`qNaKql8QaK|LJ~#|Zr9&F|m+0Gju`23KqnsM*+B>1{(A zYE-yAdUi()S?>I3Au8$y{QJagi>v8pQG{}b;AxT0H@dFGbCv&g_avVpRPs7@r{tU5K11(qecQVUShi8JArYE^M&I zS5%z8-QgPXP|{!B4VDrYDdmJ9TS}JMRd)eJ&Mn_4;cP^4aYS?R@9^S#O6)IrlWoWm z{U<(LzAk?^07WKiU%c+Vmo-FBR0yy9DzoHSz_qSe*U?r_@m-$qAy*XG6>j&|nYZYp z$b_d7f;wC-TcO_rF_mTRDBa>Usu-BTIbp|v0P&~0O3jig}W?_ zH&ClK|2+avBcQrhPLn`zfNj8r86+C4ylM*_*)%F_bhu|SS?Q*ye^Lg};(e4#U04L^ zc5&Fd%lft96TzqKyyzSlsscM)F0rpPFt+ILxzMucz24r@xz-j~mLeM70oWSsPz`Bz z=o&Byu}UMMZWP*UBn=%4mlAGAmx1UjPdSLb0*>Y*)_0UY;(7p&92PUb-NEz{(3A>Z z4JF^_nC+>?oZt((d3|Cq{fIYFiT!{6s1peNyjK zO2LGv`jzK|V?x**Y%f%us|$hpDses`I!&Xr`QoTS1t z078C0Z%YJ(gfD-p8D&EnN&5gWsv1odIUNU)gIUT$D^(pA7V1~60=4FoH%p~USVKyx zQ*6d!JF2E7SDAa}iT%@00f`Dkwb;Nu8h&W`(krUfFbg+Ba z*8XwJAB=piI4w@oP z>u!MS0z@dCNM3p<{l><3|yMa^|wRvhSagYl*$FWl~U?c%NIz^VH>?G-^@Mi)hzf9z%ze(lp?!*{ z_>mA%7=i4fyhlKZR%A;r+3LdfRys&y6I*8&he$|z6Yw{RTil}ekZH#P$&M` zS%II^jX`hex2fcuBk z>gU4}jAfPt^42Bk7gP*KqvS}`^LCX}_jF2=6R1%|j__<~RuZkaG~Hd4A@Q>#HXs8W z{(H?>T7n6X$*ois2zd40mF&=&Rm;#tZ>Y#m&P1vedXu)PceSf&{flKdzXs2pK1#DS zkvZnlaXEu0;{M=TMqg7uqmDDDT(Qqqyp&*K@GI-9rvFPehTWt3_g?w6&&0piUng(#2-6TU}NQ}SD_l$_o7B9poQ zM`Q7f;H(SNJRJ-qSdV6ev)9&*Q|LN?8(-**fEBTb3lnn&!tvr;a(#X7{ISt767-!h zh=XpQK4tbr2vBykTXr^94P3WObN6_agFT&{yHk$`z})ya`Nk$IzjruMNh36#Vp4HZ zlHi@=2JYzVS_V2kOo1e?~yrrnl6|1hB$FGB=NjaUoNjS_;-;db{4JwGiDs}R<8_$PP zRydQ|EuW$+`Fdwsf|rs1H2&A;-JW)yJ2iH37-DNwm+wo}s3VQx?mj=R5n9+e{n~6# z&mWPpm~{VcUyYL1fRKAkA6F!5ogyfH9g5T(>yvavv_Ub7*^$-x0`EhrYq6)ps@h2y zV=K>Yr?!y_>D zt?rebD=8nj;FFw_EBdG(b#WpJP*RXQGj(=_G+!nFlt;p1c=Is#(5plFrIH@iR}G7O zCU2+!0qtDyH~@PDz6mAEA|800r~h9s_yxQmbj&5hxgh0#v`Ib&oYRh%Pacy3O<_S! z0w6AYa))MGLfwU7q}4kqtnV57J+TEqWOSiiXw2I87Qrr2(mJ8N*BCLB+M1aXP$?rvZ>iRmH=Kw;lPQsl{X2 z=>*@98VIKGeHSOY1K#xh<=VtxX3cxmPnBmLpfHttD@_2v8KO$Sk7a5ZCDV+sY=D2@ zZ)m*Ymh>rvGwQH>pn6yjVdoICUW5K|)P9g4_IoA^0dqZyEF$qZ#!)DJ^Q1ik^s5Z< z3O(`cScszr)pxiR@cN#ac1ck)54}Y~gFj_sdiqywb{xhbkAtsBUg3qXU}Bf7vkD}q zzqzuThf#bZlm?0?y%S^VTR1e?ONl)QEP56*u zN>gw1-l8KC$xn;D@p*kp>_F?U+k(nJ_2z27;+kQF=KZbd^hVdYaqY^;&2!H5li;vJ z+`&61d!g~QEU4NTN7R~3L}eG(hJCF)WbiG7$tuJ!yN1uL_A_*BtF-v2pt1PAs_n$T z%vBKps)VV()@?j6(}>#>ujn917Y&~t*t_--EovO7(^p1hu`f_-%56NK@-{Om=}gPO z+p-C&eiDyqib4**FQ04&MYF<#pwZW;e55JhJI>@$@Nz87`NpGN>Xq%h&elzF{6 z-b`6<%T$ooWMw=BVQL^05mBl&=j>S!kX@Xwtu_0&FNEpfo07rbJLO<})va8J6pN=a zMX);f7RmBircQWIiru}0`BPLB#@USGeoTsO>==M{0C?443DBU!U7zK^VJsrLhKp)Z zgWQbV)$BGq{zDtHOe(7QU6GRW)r6=xV&X4&@$^fPG7{A9{4%)u-k!Rw(x=|m}=EBKs>RMA8<0XEQvs?lYVJB1mYAWXbw@9#Qe#484 zC%NcJLi?q~m-6JF^J^*^J-oDU_A`}#K1mzo@`aHZ$xn3L&l?o`=t@-yW>sZcN4Ja< z+p&OEaq3$o*_jC99SI75Z?c2&6amY@OAo{@c}d&pi9vIpcC52c4`j=Zze06ltQ?mp z*-=6uX}_ADV6xDo_2*Dl1eSzyfygCtKo<0#&-k`K6Lcrb+|L`vT!CYb3Y!Ad;P@B; zAY?%dB#GltZ8rvKmzplFG~JfLV;*>mTqk5^AqPmKch&MW1K*E~=8ai~zXIcBllRAX zS=A>oUL9Tf;$g!k;+NwaQfeERU_Bu#Y6HnFll|r{+bvZ)E$LaubFkrqcCGD`qq?FF zCpY`n^w5J&JeBDJ6`kN_R1)^Pb`dWT$4^s{r0okoLp7X4Cc zTJ-FR=ZxRBuM2XX{^kTZfjzTX&5}nMMqu^Tk)+4%)C zQVD$eK!_95ll?XsCxrInDx%exARaZlwk|8H1tj;dwlll!|@aiBXA<--ZA6e%9dqD6LdJ z%?>9&DIii_TX+rfi$8R~f2`tpu!=KAT+We9tt&?G;nd<6_O{@E?mh$tzI^7PA|BYh zV0<6_5(rpXyTXNk+kjnwrMwb-m;y{1vtrwq=xCo%#MOnj#=qd7faB>hedpZ^WKacU zw7Jdw(cc4xzm~-N?*463wn(>vNxB zd}#Pi@ek_-e&~tOcj_gQKY4jz<``#<2jo9o{V$vUgfoI7%DN<7o^X%_2OWQ3Q$43h ziY0rqWfKzh!;8F;F5j)YLW8izYTZk{%KyznP6yk?Z5d#zvr!*6K_nB-5sDv)V@ag z6X5`$Lx1`v(6S*(m)!YU)yV|wpm09+jxm+7=3mgkz`*5Z@FDDhW-49Bc5r)PafKgD zTMz-@Hh(6aCa!Zh)dYh(Ip6b5hn>RPbSU^UE6!&C#RR?{Y%#~Qg>{bg5;-)idhgGu z-F|?3G-Ah?OjORih2GA}h>e26DeOPvC@$XDGZP3`nRgN4+Edh^eG|4}Gj1hJ2ni>l$$F7`R83VbJ^;_>ss@pi?O)xUSZ z=|0ER1GnbUTXTW1EuO-UJ+C_Yj-g)L>YDe%m+a9*v6YubVD>S_ zfX`q~)%0pp)yeuQJ5w+}X2SUK4%nw_YDT z#iR!TS0!a4l5{tDN@5T*m~hXz29vBY(T&}*cA4+?_8tM@Oi1k*O-#J8H@a7~-21RX zhULfCnCZdSB7;P5x1fsbp}_rn_U4aYElzaee*1Xt3hvuJ0ssHQ|2v(~luz9WkBqE9 zB9UG|h4yUm(~>cv@87=z@vZS5eB97fv8?p09)4e(0D$ar7!2lx93Ni=@X~;)+%;$A z`T|fE?Ztu{%e@tySI|&4nMwNRogE#|n3$MutrSg!l)3LebpdRRzyH&SYBvBcFb|bK z1x&l{huiZ#s4lM%FbN?(7>V7^l7L;)%iY|;E5o__-arTsy%bsUbYgide^@BG+70Nt zS1+%AtOhzI>npi6tUsn%-0h93eZ&1a5;lr#%(cBl-Hawh$Plp^_=@f67~GvV8=hsQ zJyyAQ&j~OJvgjFhX09!s!L?gF(A3}wJz#p#PaJ6sW+Y@Ev1qGl*X{_Ujg7_;w_UHN|t+@XBJ2=tPfAg zljdzJExh8W)5#HQ2RB5MA!b0q1;!O0@+KIpvFX69+}wSb4yng7#KOiVJNHveZ0t&A z290dWSw5xo0!yAo#d$4M>q?NrL!Lwp`>*vcrFny&XPBb0y47CfF>4eLE@psw*Q}b{ z_y~8#!8=GpIL)^~6GN02sDSNxK9RFQ3Yq>)3Kn;W?2G)n97CTi5t>?_K$^o-ODlbd z8>USfT9B=20Bg(KiB1n94Ogty%0TKZp^q_WPtKz7xXYet=lK#fyP_TE^xEA73@gPQ zoOEWH9&i~3VV#kT*HGj$>F-1 z5U_XjfqA&IlKi0+md8nFT;#7%RIZnqQ?7VdJC!>%W)0LmW1jPQMjCVHve9|S<#2t~ zOWz(H8Tpo56eosWcHFpQ=nZso69;z(os%=Z7y|Timd%3~Ma~`&NRz^khp2PGNaiCq zxu~|4bCkoB4>TWG(UG>&z+oq!me_HjL7vyeBS+6|PnR$8;Wb-$gDbIK7@BkTz3lpE zkyXQoVMW;J=VcKv$TYY?E>O0Qi;5bdm*j|2>~T72jxw#!e{T5*+|U(fS*(p=&X@G& z9qw0@g{`Ig^9XY|r1@mSt;UuZO-Y+r0hxAT`FfoniGHFHN@Z7{F%tEPQ23>YX9(Gn zI;^<=*?ovlGCMH0;%2)89Z8bGS&}Q(0@tMR)VT`)t5&gAj2ue!oimW(F==;38fA>D z_d0!F=rrkep_$?M!reSSO4DL)X@-QSNPez?+htR2Tc#*xRZ8!B0!~LS{qyGg=8nooNe1hITxf^}cc1X?Hv4vrX4jw=)SU=QdP2kek{|)-{}e{t zqPwE`V}Z-w0>c3=%EX_4^@rkq>pA)V(ua2VSvfh*z?SGM0W7EO7~C{ywOu_nAz=e& zO@~@H0kh2UBq$yW;LB}JdOF;;Rn|SeQNg<(jplQx*)27j!rA1}1hs@%zm`J>gOt$n zc&oAWjfuUgc84j(1n1Xp3^2#|klD*u2z*TW@p=G&xB_Re%j{>Hw8b4jd6U=J%*-S2 zjtUE$=>=+5)%gBEPHJZsx6ao=iG9P6OFs~_++Rf^5(k|GjH*{f#l&7|4{u!eMD6rQ zJa_ZHbonxR3vdph_D6AM85>eCR$Uc0=W}I0eya@ddSJbvQtfK3SSF`ZD#)q|7?tL6 zP*H>y7?i_L^g!lt@G6$>^&$I{P6Odm)9h2D-@}W%B9b|))a*e$A95Y zBYebhkq1vpH*_`g7UGS9d(EMKM!e!l{kenqrj&J_*;q$@Uf!L>j^b`3V|9M!Npk4=7=@j?rOk@VH2O`l{-ar*b#{aIqIgI%Ex6><>&depHpCBhEmW7F0g(VgF)BMz}9rVx93{5*%)Jipt`5+oBdgk>))t8m3(#yju`3jpmU<0Zi<-LIfn-Hq|-kx03IdmHa%BzuInxEgBottZWo!`OwW~ths z(X~r>HeILYC3g?x$M`+(c=D;$s9xVqSPNy9T9%nD+xK_Swe9{ce~e)(UzhO~z0x#I zjjdieo-=yYw8GWA_E=&+eqVUZVZ*$mz!S*wu2d;h9AWQ=%3%&UX_8>wB#=DNdPtt? zELq2+nnlIHcaV#|y7(!{!d}6NV|)}GFPt_X?enJQnq;ua3=ZI%I-iHpf=iBAvLxWlR&+ zPS@sM^LM&7a%)du*WgxZsFYN7{&Qv@wXC2H{%Goom}6M#SzMt0Z??&vN6{= zT>}%1Py@JFcD2IpaB1OUTUBIRh2}{>TT^W09lIE9{*Wdv#WMYKGneqtZgIo$DvJ|# zD!eV%{%rrDx;eb9-9K~}O$u2JU5bX-b(k+}AD;8M3yBEIR9CoSNoG*ZKHPtGY)>J~ zJv%K5mVY-zGrE}GT+awU`T28b#9F#CcmGtM>WjqZ%cBaKsk3rrQ<{_W2sPD4zHT0# z{e+Vou-D*-wLUg4qS+(@yO-dFT58#Pg63{3ideB2OW9Bsd4t;BMVD$|cZJkaIWFDL z>2{=9yCUT{P}u(5{nLA1EGrEPx=O=cVyYHe1=^VB4GivdR6E1CCN^Tp(haV`#92rK zpKB&)vkG~RS|~N+rSYO8z28PVC`^ffrCXC;Tn2l0xHDq{?<>(k+c(?NCHwFtI8s}t zf({kn z3>FSJ6npQ@6HOm$N{#w}S4$03{OKTk{y*y(jHIp{Q`>8*uXX3p(cv}+1v!=5IX8`8 F{1=-lqX_^2 literal 0 HcmV?d00001 diff --git a/src/main.cpp b/src/main.cpp index 6779123..fd74aea 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,7 +14,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 25; // feel free to change the size of array +const int SIZE = 1 << 27; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; @@ -138,13 +138,27 @@ int main(int argc, char* argv[]) { printDesc("work-efficient compact, power-of-two"); count = StreamCompaction::Efficient::compact(SIZE, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(count, c, true); + printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); printDesc("work-efficient compact, non-power-of-two"); count = StreamCompaction::Efficient::compact(NPOT, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + printArray(count, c, true); + printCmpLenResult(count, expectedNPOT, b, c); + + zeroArray(SIZE, c); + printDesc("thrust compact, power-of-two"); + count = StreamCompaction::Thrust::compact(SIZE, c, a); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(count, c, true); + printCmpLenResult(count, expectedCount, b, c); + + zeroArray(SIZE, c); + printDesc("thrust compact, non-power-of-two"); + count = StreamCompaction::Thrust::compact(NPOT, c, a); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); @@ -167,6 +181,7 @@ int main(int argc, char* argv[]) { std::memcpy(b, a, sizeof(int) * SIZE); std::sort(b, b + SIZE); StreamCompaction::Efficient::radix_sort(SIZE, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); printCmpResult(SIZE, b, c); printArray(SIZE, c, true); @@ -174,6 +189,7 @@ int main(int argc, char* argv[]) { std::memcpy(b, a, sizeof(int) * NPOT); std::sort(b, b + NPOT); StreamCompaction::Efficient::radix_sort(NPOT, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); printCmpResult(NPOT, b, c); printArray(NPOT, c, true); diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 3ede5e4..8558538 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -41,8 +41,7 @@ namespace StreamCompaction { return; } - int val = bools[iSelf]; - if (val != 0) { + if (bools[iSelf] != 0) { odata[indices[iSelf]] = idata[iSelf]; } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index c69e76a..3c5da31 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -15,36 +15,50 @@ namespace StreamCompaction { return timer; } - // block size when processing. cuda block size is half this - constexpr int log2BlockSize = 9; // 512 + // this many elements are processed by one block. cuda block size is half this + /*constexpr int log2BlockSize = 10; // 1024*/ + /*constexpr int log2BlockSize = 9; // 512*/ + /*constexpr int log2BlockSize = 8; // 256*/ + constexpr int log2BlockSize = 7; // 128 + /*constexpr int log2BlockSize = 6; // 64*/ constexpr int blockSize = 1 << log2BlockSize; + constexpr int log2BankSize = 5; + + __device__ int conflictFreeIndex(int i) { + return i + (i >> log2BankSize); + } + __global__ void kernScanPerBlock(int *data, int *lastData) { extern __shared__ int buffer[]; data += blockIdx.x * blockDim.x * 2; + int offset1 = conflictFreeIndex(threadIdx.x), offset2 = conflictFreeIndex(threadIdx.x + blockDim.x); + // copy data to shared memory - buffer[threadIdx.x * 2] = data[threadIdx.x * 2]; - buffer[threadIdx.x * 2 + 1] = data[threadIdx.x * 2 + 1]; - __syncthreads(); + buffer[offset1] = data[threadIdx.x]; + buffer[offset2] = data[threadIdx.x + blockDim.x]; int lastElem = 0; if (lastData && threadIdx.x == blockDim.x - 1) { - lastElem = buffer[threadIdx.x * 2 + 1]; + lastElem = buffer[offset2]; } + __syncthreads(); // upward pass for (int halfGap = 1; halfGap < blockDim.x; halfGap <<= 1) { if (threadIdx.x < blockDim.x / halfGap) { - buffer[(threadIdx.x * 2 + 2) * halfGap - 1] += buffer[(threadIdx.x * 2 + 1) * halfGap - 1]; + int + id1 = conflictFreeIndex((threadIdx.x * 2 + 1) * halfGap - 1), + id2 = conflictFreeIndex((threadIdx.x * 2 + 2) * halfGap - 1); + buffer[id2] += buffer[id1]; } __syncthreads(); } if (threadIdx.x == blockDim.x - 1) { - int halfIdx = blockDim.x - 1; - buffer[blockDim.x * 2 - 1] = buffer[threadIdx.x]; - buffer[threadIdx.x] = 0; + buffer[conflictFreeIndex(blockDim.x * 2 - 1)] = buffer[offset1]; + buffer[offset1] = 0; } __syncthreads(); @@ -53,6 +67,8 @@ namespace StreamCompaction { if (threadIdx.x < blockDim.x / halfGap) { int prevIdx = (threadIdx.x * 2 + 1) * halfGap - 1; int thisIdx = prevIdx + halfGap; + prevIdx = conflictFreeIndex(prevIdx); + thisIdx = conflictFreeIndex(thisIdx); int sum = buffer[thisIdx] + buffer[prevIdx]; buffer[prevIdx] = buffer[thisIdx]; buffer[thisIdx] = sum; @@ -61,10 +77,10 @@ namespace StreamCompaction { } // copy data back - data[threadIdx.x * 2] = buffer[threadIdx.x * 2]; - data[threadIdx.x * 2 + 1] = buffer[threadIdx.x * 2 + 1]; + data[threadIdx.x] = buffer[offset1]; + data[threadIdx.x + blockDim.x] = buffer[offset2]; if (lastData && threadIdx.x == blockDim.x - 1) { - lastData[blockIdx.x] = lastElem + buffer[threadIdx.x * 2 + 1]; + lastData[blockIdx.x] = lastElem + buffer[offset2]; } } @@ -90,13 +106,17 @@ namespace StreamCompaction { int *buffer; cudaMalloc(&buffer, sizeof(int) * indirectSize); - kernScanPerBlock<<>>(dev_data, buffer); + kernScanPerBlock<<< + numBlocks, blockSize / 2, (blockSize + (blockSize >> log2BankSize)) * sizeof(int) + >>>(dev_data, buffer); dev_scan(indirectSize, buffer); kernAddConstantToBlock<<>>(dev_data, buffer); cudaFree(buffer); } else { // just scan the block - kernScanPerBlock<<<1, blockSize / 2, blockSize * sizeof(int)>>>(dev_data, nullptr); + kernScanPerBlock<<< + 1, blockSize / 2, (blockSize + (blockSize >> log2BankSize)) * sizeof(int) + >>>(dev_data, nullptr); } } @@ -137,10 +157,12 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - constexpr int log2BlockSize = 9; // block_size = 512 + constexpr int log2ScatterBlockSize = 6; + constexpr int scatterBlockSize = 1 << log2ScatterBlockSize; int numBlocks, bufferSize; _computeSizes(n, log2BlockSize, &numBlocks, &bufferSize); + int numScatterBlocks = (n + scatterBlockSize - 1) >> log2ScatterBlockSize; int *data, *accum, *out; cudaMalloc(&data, sizeof(int) * bufferSize); @@ -148,21 +170,23 @@ namespace StreamCompaction { cudaMalloc(&out, sizeof(int) * bufferSize); cudaMemcpy(data, idata, sizeof(int) * n, cudaMemcpyHostToDevice); - cudaMemset(data + n, 0, sizeof(int) * (bufferSize - n)); timer().startGpuTimer(); Common::kernMapToBoolean<<>>(n, accum, data); dev_scan(bufferSize, accum); - Common::kernScatter<<>>(n, out, data, data, accum); + Common::kernScatter<<>>(n, out, data, data, accum); timer().endGpuTimer(); - int res; - cudaMemcpy(odata, out, sizeof(int) * n, cudaMemcpyDeviceToHost); + int last = idata[n - 1] != 0 ? 1 : 0, res; cudaMemcpy(&res, accum + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + res += last; + cudaMemcpy(odata, out, sizeof(int) * res, cudaMemcpyDeviceToHost); checkCUDAError("efficient compaction"); - if (idata[n - 1] != 0) { - ++res; - } + + cudaFree(data); + cudaFree(accum); + cudaFree(out); + return res; } @@ -210,6 +234,7 @@ namespace StreamCompaction { cudaMemcpy(data1, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + timer().startGpuTimer(); for (int i = 0; i < numIntBits; ++i) { kernExtractBit<<>>(n, i, trues, data1); kernNegate<<>>(falses, trues); @@ -224,6 +249,7 @@ namespace StreamCompaction { kernRadixSortScatter<<>>(n, numFalses, i, data2, data1, trues, falses); std::swap(data1, data2); } + timer().endGpuTimer(); cudaMemcpy(odata, data1, sizeof(int) * n, cudaMemcpyDeviceToHost); diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index e003bf5..3286b33 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -12,7 +12,8 @@ namespace StreamCompaction { return timer; } - constexpr int block_size = 128; + constexpr int log2BlockSize = 7; + constexpr int block_size = 1 << log2BlockSize; __global__ void kernScanPass(int n, int diff, int *odata, const int *idata) { int iSelf = blockIdx.x * blockDim.x + threadIdx.x; diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 0ee6ce5..2769bf5 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -29,5 +29,24 @@ namespace StreamCompaction { thrust::copy(dev_out.begin(), dev_out.end(), odata); } + + struct isNonZero { + __host__ __device__ bool operator()(int x) { + return x != 0; + } + }; + + int compact(int n, int *odata, const int *idata) { + thrust::device_vector dev_vec(idata, idata + n), dev_out(n); + + timer().startGpuTimer(); + int count = thrust::copy_if( + dev_vec.begin(), dev_vec.end(), dev_out.begin(), isNonZero() + ) - dev_out.begin(); + timer().endGpuTimer(); + + thrust::copy(dev_out.begin(), dev_out.end(), odata); + return count; + } } } diff --git a/stream_compaction/thrust.h b/stream_compaction/thrust.h index fe98206..a8991f9 100644 --- a/stream_compaction/thrust.h +++ b/stream_compaction/thrust.h @@ -7,5 +7,7 @@ namespace StreamCompaction { StreamCompaction::Common::PerformanceTimer& timer(); void scan(int n, int *odata, const int *idata); + + int compact(int n, int *odata, const int *idata); } }