From 4d5da87badd81ccdb601ed9509372995522ac306 Mon Sep 17 00:00:00 2001 From: liamdugan Date: Wed, 19 Sep 2018 07:21:29 -0400 Subject: [PATCH 1/4] temp commit to go work in moore --- src/main.cpp | 10 ++-- src/testing_helpers.hpp | 14 +++--- stream_compaction/CMakeLists.txt | 2 +- stream_compaction/common.cu | 24 +++++++++- stream_compaction/cpu.cu | 62 +++++++++++++++++++----- stream_compaction/efficient.cu | 10 ++-- stream_compaction/naive.cu | 82 ++++++++++++++++++++++++++++---- stream_compaction/naive.h | 1 + stream_compaction/thrust.cu | 10 ++-- 9 files changed, 170 insertions(+), 45 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 1850161..8fc5ab0 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -60,12 +60,12 @@ int main(int argc, char* argv[]) { StreamCompaction::Naive::scan(SIZE, c, a); printArray(SIZE, c, true); */ - zeroArray(SIZE, c); - printDesc("naive scan, non-power-of-two"); - StreamCompaction::Naive::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + // zeroArray(SIZE, c); + // printDesc("naive scan, non-power-of-two"); + // StreamCompaction::Naive::scan(NPOT, c, a); + // printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); //printArray(SIZE, c, true); - printCmpResult(NPOT, b, c); + // printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, power-of-two"); diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index 46337ab..d6cc4e3 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -1,8 +1,8 @@ #pragma once -#include -#include -#include +#include +#include +#include #include #include @@ -69,8 +69,8 @@ void printArray(int n, int *a, bool abridged = false) { printf("]\n"); } -template -void printElapsedTime(T time, std::string note = "") -{ - std::cout << " elapsed time: " << time << "ms " << note << std::endl; +template +void printElapsedTime(T time, std::string note = "") +{ + std::cout << " elapsed time: " << time << "ms " << note << std::endl; } \ No newline at end of file diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index cdbef77..48e2f35 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -13,5 +13,5 @@ set(SOURCE_FILES cuda_add_library(stream_compaction ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_60 ) diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 8fc0211..659e7b6 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -23,7 +23,16 @@ 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 + + // get index first + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) + { + return; + } + + // determine if you're a boolean + bools[index] = idata[index] ? 1 : 0; } /** @@ -32,7 +41,18 @@ namespace StreamCompaction { */ __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { - // TODO + + // get index first + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) + { + return; + } + + if (bools[index]) + { + odata[indices[index]] = idata[index]; + } } } diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 05ce667..a5bd5db 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -1,15 +1,15 @@ #include #include "cpu.h" -#include "common.h" +#include "common.h" namespace StreamCompaction { namespace CPU { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } /** @@ -19,7 +19,14 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + + int sum = 0; + for (int i = 0; i < n; ++i) + { + odata[i] = sum; + sum += idata[i]; + } + timer().endCpuTimer(); } @@ -30,9 +37,20 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + + int index = 0; + for (int i = 0; i < n; ++i) + { + // if the data meets the condition put it in + if (idata[i]) + { + odata[index] = idata[i]; + ++index; + } + } + timer().endCpuTimer(); - return -1; + return index; } /** @@ -42,9 +60,31 @@ namespace StreamCompaction { */ int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + + // then run the scan on it + int* scanned = (int*) malloc(sizeof(int) * n); + int sum = 0; + for (int i = 0; i < n; ++i) + { + scanned[i] = sum; + if (idata[i]) + { + ++sum; + } + } + + // now scatter + for (int j = 0; j < n; j++) + { + if (idata[j]) + { + odata[scanned[j]] = idata[j]; + } + } + + free(scanned); timer().endCpuTimer(); - return -1; + return sum; } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 36c5ef2..fd1d622 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -5,11 +5,11 @@ namespace StreamCompaction { namespace Efficient { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } /** diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 9218f8e..779312d 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -3,23 +3,87 @@ #include "common.h" #include "naive.h" +#define blockSize 512 + + __global__ void kernNaiveScan(int n, int twoToPowerDMinusOne, float* odata, float* idata) + { + // get index first + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) + { + return; + } + + // then add the two numbers and put them into the global output buffer + // if (index >= twoToPowerDMinusOne) + // { + // odata[index] = idata[index - twoToPowerDMinusOne] + idata[index]; + // } + // else + // { + // odata[index] = idata[index]; + // } + } + namespace StreamCompaction { namespace Naive { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } - // TODO: __global__ + /** * 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(); + timer().startGpuTimer(); + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + + float* dev_gpuScanBuf; + float* dev_idata; + + int nNextHighestPowTwo = 1 << ilog2ceil(n); + + cudaMalloc((void**)&dev_gpuScanBuf, nNextHighestPowTwo * sizeof(float)); + checkCUDAError("cudaMalloc buf failed"); + + cudaMalloc((void**)&dev_idata, nNextHighestPowTwo * sizeof(float)); + checkCUDAError("cudaMalloc idata failed"); + + cudaMemcpy((void*)dev_idata, (const void*)idata, nNextHighestPowTwo * sizeof(float), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // call the kernel log2n number of times + bool flipped = false; + for (int i = 0; i < ilog2ceil(nNextHighestPowTwo); ++i) + { + // call the kernel + int twoToPowerIMinusOne = 1 << (i - 1); + std::cout << ((n + blockSize - 1) / blockSize) << ", " << blockSize << std::endl; + kernNaiveScan<<<((n + blockSize - 1) / blockSize) , blockSize>>>(nNextHighestPowTwo, twoToPowerIMinusOne, dev_gpuScanBuf, dev_idata); + + // flip flop the buffers and keep track with a boolean (flipped = true means dev_idata has the latest data) + float* temp = dev_gpuScanBuf; + dev_gpuScanBuf = dev_idata; + dev_idata = temp; + flipped = !flipped; + } + + if (flipped) + { + cudaMemcpy(odata, dev_idata, nNextHighestPowTwo * sizeof(float), cudaMemcpyDeviceToHost); + } + else + { + cudaMemcpy(odata, dev_gpuScanBuf, nNextHighestPowTwo * sizeof(float), cudaMemcpyDeviceToHost); + } + + cudaFree(dev_gpuScanBuf); + cudaFree(dev_idata); + timer().endGpuTimer(); } } } diff --git a/stream_compaction/naive.h b/stream_compaction/naive.h index 37dcb06..bf2d3b5 100644 --- a/stream_compaction/naive.h +++ b/stream_compaction/naive.h @@ -1,6 +1,7 @@ #pragma once #include "common.h" +#include namespace StreamCompaction { namespace Naive { diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 36b732d..e3b3268 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -8,11 +8,11 @@ namespace StreamCompaction { namespace Thrust { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. From 1b51b3bbace328a287342c38c2b24be6da4f92f2 Mon Sep 17 00:00:00 2001 From: liamdugan Date: Thu, 20 Sep 2018 21:45:03 -0400 Subject: [PATCH 2/4] Finished code --- src/main.cpp | 57 +++++----- stream_compaction/common.cu | 14 ++- stream_compaction/common.h | 2 +- stream_compaction/efficient.cu | 188 ++++++++++++++++++++++++++++++++- stream_compaction/naive.cu | 92 +++++++++------- stream_compaction/thrust.cu | 11 +- 6 files changed, 280 insertions(+), 84 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 8fc5ab0..3adb73e 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,11 +13,12 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array -const int NPOT = SIZE - 3; // Non-Power-Of-Two +const int SIZE = 1 << 20; // feel free to change the size of array +const int NPOT = SIZE - 30; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; int *c = new int[SIZE]; +int *d = new int[SIZE]; int main(int argc, char* argv[]) { // Scan tests @@ -29,7 +30,7 @@ int main(int argc, char* argv[]) { genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; - printArray(SIZE, a, true); +// printArray(SIZE, a, true); // initialize b using StreamCompaction::CPU::scan you implement // We use b for further comparison. Make sure your StreamCompaction::CPU::scan is correct. @@ -38,61 +39,61 @@ int main(int argc, char* argv[]) { printDesc("cpu scan, power-of-two"); StreamCompaction::CPU::scan(SIZE, b, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(SIZE, b, true); + // printArray(SIZE, b, true); zeroArray(SIZE, c); printDesc("cpu scan, non-power-of-two"); StreamCompaction::CPU::scan(NPOT, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(NPOT, b, true); +// printArray(NPOT, b, true); printCmpResult(NPOT, b, c); - + zeroArray(SIZE, c); printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); +// printArray(SIZE, c, true); printCmpResult(SIZE, b, c); - /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan - onesArray(SIZE, c); - printDesc("1s array for finding bugs"); - StreamCompaction::Naive::scan(SIZE, c, a); - printArray(SIZE, c, true); */ + /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan */ + onesArray(SIZE, d); + printDesc("1s array for finding bugs"); + StreamCompaction::Efficient::scan(SIZE, c, d); +// printArray(SIZE, c, true); - // zeroArray(SIZE, c); - // printDesc("naive scan, non-power-of-two"); - // StreamCompaction::Naive::scan(NPOT, c, a); - // printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); - // printCmpResult(NPOT, b, c); + zeroArray(SIZE, c); + printDesc("naive scan, non-power-of-two"); + StreamCompaction::Naive::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); +// printArray(SIZE, c, true); + printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, power-of-two"); StreamCompaction::Efficient::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); +// printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, non-power-of-two"); StreamCompaction::Efficient::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(NPOT, c, true); +// printArray(NPOT, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("thrust scan, power-of-two"); StreamCompaction::Thrust::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); +// printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("thrust scan, non-power-of-two"); StreamCompaction::Thrust::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(NPOT, c, true); +// printArray(NPOT, c, true); printCmpResult(NPOT, b, c); printf("\n"); @@ -104,7 +105,7 @@ int main(int argc, char* argv[]) { genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; - printArray(SIZE, a, true); +// printArray(SIZE, a, true); int count, expectedCount, expectedNPOT; @@ -115,7 +116,7 @@ int main(int argc, char* argv[]) { count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); expectedCount = count; - printArray(count, b, true); +// printArray(count, b, true); printCmpLenResult(count, expectedCount, b, b); zeroArray(SIZE, c); @@ -123,28 +124,28 @@ int main(int argc, char* argv[]) { count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); expectedNPOT = count; - printArray(count, c, true); + // printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); zeroArray(SIZE, c); printDesc("cpu compact with scan"); count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(count, c, true); + // printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); 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); +// printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); system("pause"); // stop Win32 console from closing on exit diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 659e7b6..25f9fc3 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -22,17 +22,16 @@ namespace StreamCompaction { * Maps an array to an array of 0s and 1s for stream compaction. Elements * 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) { - - // get index first + __global__ void kernMapToBoolean(int n, int paddedN, int *bools, const int *idata) { + // get index first and reject if greater than paddedN int index = threadIdx.x + (blockIdx.x * blockDim.x); - if (index >= n) + if (index >= paddedN) { return; } - - // determine if you're a boolean - bools[index] = idata[index] ? 1 : 0; + + // determine if you're a boolean (if you're in the part that's just padded on, give yourself a 0) + bools[index] = (idata[index] && index < n) ? 1 : 0; } /** @@ -54,6 +53,5 @@ namespace StreamCompaction { odata[indices[index]] = idata[index]; } } - } } diff --git a/stream_compaction/common.h b/stream_compaction/common.h index 99a1b04..13a4059 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -32,7 +32,7 @@ inline int ilog2ceil(int x) { namespace StreamCompaction { namespace Common { - __global__ void kernMapToBoolean(int n, int *bools, const int *idata); + __global__ void kernMapToBoolean(int n, int paddedN, int *bools, const int *idata); __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices); diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index fd1d622..80c2ea7 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -1,8 +1,54 @@ + #include #include #include "common.h" #include "efficient.h" +#define blockSize 512 + +int* dev_efficientScanBuf; +int* dev_efficientIdata; +int* dev_efficientBools; +int* dev_efficientIndices; + +__global__ void kernEfficientScanUpSweep(int n, int d, int* odata, int* idata) +{ + // get index first + int index = threadIdx.x + (blockIdx.x * blockDim.x); + int twoToPowDPlusOne = 1 << (d + 1); + if (index >= n || index % twoToPowDPlusOne != 0) + { + return; + } + + int twoToPowD = 1 << d; + + // then add the two numbers and put them into the global output buffer + odata[index + twoToPowDPlusOne - 1] = idata[index + twoToPowDPlusOne - 1] + idata[index + twoToPowD - 1]; +} + +__global__ void kernSetFirstElementZero(int n, int* odata) +{ + odata[n - 1] = 0; +} + +__global__ void kernEfficientScanDownSweep(int n, int d, int* odata, int* idata) +{ + // get index first + int index = threadIdx.x + (blockIdx.x * blockDim.x); + int twoToPowDPlusOne = 1 << (d + 1); + if (index >= n || (index % twoToPowDPlusOne != 0)) + { + return; + } + + int twoToPowD = 1 << d; + + // then sweep down + odata[index + twoToPowD - 1] = idata[index + twoToPowDPlusOne - 1]; + odata[index + twoToPowDPlusOne - 1] = idata[index + twoToPowDPlusOne - 1] + idata[index + twoToPowD - 1]; +} + namespace StreamCompaction { namespace Efficient { using StreamCompaction::Common::PerformanceTimer; @@ -16,9 +62,61 @@ namespace StreamCompaction { * 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(); + timer().startGpuTimer(); + int nNextHighestPowTwo = 1 << ilog2ceil(n); + + cudaMalloc((void**)&dev_efficientScanBuf, nNextHighestPowTwo * sizeof(int)); + checkCUDAError("cudaMalloc buf failed"); + + cudaMalloc((void**)&dev_efficientIdata, nNextHighestPowTwo * sizeof(int)); + checkCUDAError("cudaMalloc idata failed"); + + cudaMemcpy((void*)dev_efficientIdata, (const void*)idata, nNextHighestPowTwo * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // call the upsweep kernel log2n number of times + for (int d = 0; d < ilog2ceil(nNextHighestPowTwo); ++d) + { + + // copy all the data to make sure everythings in place + cudaMemcpy((void*)dev_efficientScanBuf, (const void*)dev_efficientIdata, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // call the kernel + kernEfficientScanUpSweep<<<((nNextHighestPowTwo + blockSize - 1) / blockSize) , blockSize>>>(nNextHighestPowTwo, d, dev_efficientScanBuf, dev_efficientIdata); + + // flip flop the buffers so that idata is always the most recent data + int* temp = dev_efficientScanBuf; + dev_efficientScanBuf = dev_efficientIdata; + dev_efficientIdata = temp; + } + + // set first element to be zero in a new kernel (unsure how to do this otherwise) + kernSetFirstElementZero << <1, 1 >> > (nNextHighestPowTwo, dev_efficientIdata); + + // now call the downsweep kernel log2n times + for (int d = (ilog2ceil(nNextHighestPowTwo) - 1); d >= 0; --d) + { + // copy all the data to make sure everythings in place + cudaMemcpy((void*)dev_efficientScanBuf, (const void*)dev_efficientIdata, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // call the kernel + kernEfficientScanDownSweep<<<((nNextHighestPowTwo + blockSize - 1) / blockSize) , blockSize>>>(nNextHighestPowTwo, d, dev_efficientScanBuf, dev_efficientIdata); + + // flip flop the buffers + int* temp = dev_efficientScanBuf; + dev_efficientScanBuf = dev_efficientIdata; + dev_efficientIdata = temp; + } + + // shift it and memcpy to out + cudaMemcpy(odata, dev_efficientIdata, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToHost); + + cudaFree(dev_efficientScanBuf); + cudaFree(dev_efficientIdata); + + timer().endGpuTimer(); } /** @@ -32,9 +130,89 @@ namespace StreamCompaction { */ int compact(int n, int *odata, const int *idata) { timer().startGpuTimer(); - // TODO + int nNextHighestPowTwo = 1 << ilog2ceil(n); + + cudaMalloc((void**)&dev_efficientBools, nNextHighestPowTwo * sizeof(int)); + checkCUDAError("cudaMalloc bool buf failed"); + + cudaMalloc((void**)&dev_efficientScanBuf, nNextHighestPowTwo * sizeof(int)); + checkCUDAError("cudaMalloc buf failed"); + + cudaMalloc((void**)&dev_efficientIdata, nNextHighestPowTwo * sizeof(int)); + checkCUDAError("cudaMalloc idata failed"); + + cudaMalloc((void**)&dev_efficientIndices, nNextHighestPowTwo * sizeof(int)); + checkCUDAError("cudaMalloc indices failed"); + + // memcpy all the stuff over to gpu before calling kernel functions + cudaMemcpy((void*)dev_efficientIdata, (const void*)idata, nNextHighestPowTwo * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // map all of the values to booleans (and pad with zeroes for those values higher than original array limit) + StreamCompaction::Common::kernMapToBoolean<< <((nNextHighestPowTwo + blockSize - 1) / blockSize), blockSize >> > (n, nNextHighestPowTwo, dev_efficientBools, dev_efficientIdata); + + // Start the scan --------------- (copy pasted from the scan function because you can't nest calls to timer. Plus it saves a copy from device to host) + + // make a copy of the bools so we can do the scan and put it into indices + cudaMemcpy((void*)dev_efficientIndices, (const void*)dev_efficientBools, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // call the upsweep kernel log2n number of times + for (int d = 0; d < ilog2ceil(nNextHighestPowTwo); ++d) + { + // copy all the data to make sure everythings in place + cudaMemcpy((void*)dev_efficientScanBuf, (const void*)dev_efficientIndices, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // call the kernel + kernEfficientScanUpSweep << <((nNextHighestPowTwo + blockSize - 1) / blockSize), blockSize >> > (nNextHighestPowTwo, d, dev_efficientScanBuf, dev_efficientIndices); + + // flip flop the buffers so that idata is always the most recent data + int* temp = dev_efficientScanBuf; + dev_efficientScanBuf = dev_efficientIndices; + dev_efficientIndices = temp; + } + + // set first element to be zero in a new kernel (unsure how to do this otherwise) + kernSetFirstElementZero << <1, 1 >> > (nNextHighestPowTwo, dev_efficientIndices); + + // now call the downsweep kernel log2n times + for (int d = (ilog2ceil(nNextHighestPowTwo) - 1); d >= 0; --d) + { + // copy all the data to make sure everythings in place + cudaMemcpy((void*)dev_efficientScanBuf, (const void*)dev_efficientIndices, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy idata failed"); + + // call the kernel + kernEfficientScanDownSweep << <((nNextHighestPowTwo + blockSize - 1) / blockSize), blockSize >> > (nNextHighestPowTwo, d, dev_efficientScanBuf, dev_efficientIndices); + + // flip flop the buffers + int* temp = dev_efficientScanBuf; + dev_efficientScanBuf = dev_efficientIndices; + dev_efficientIndices = temp; + } + + // ------- end of scan + + int sizeOfCompactedStream = 0; + // memcpy the final value of indices to out so that we can get the total size of compacted stream + cudaMemcpy(&sizeOfCompactedStream, dev_efficientIndices + (nNextHighestPowTwo - 1), 1 * sizeof(int), cudaMemcpyDeviceToHost); + + // run the stream compaction + StreamCompaction::Common::kernScatter << <((nNextHighestPowTwo + blockSize - 1) / blockSize), blockSize >> > (nNextHighestPowTwo, dev_efficientScanBuf, dev_efficientIdata, dev_efficientBools, dev_efficientIndices); + + // memcpy to out + cudaMemcpy(odata, dev_efficientScanBuf, sizeOfCompactedStream * sizeof(int), cudaMemcpyDeviceToHost); + + // free all our stuff + cudaFree(dev_efficientScanBuf); + cudaFree(dev_efficientBools); + cudaFree(dev_efficientIdata); + cudaFree(dev_efficientIndices); timer().endGpuTimer(); - return -1; + + // return the total size of the compacted stream + return sizeOfCompactedStream; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 779312d..576738e 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -5,25 +5,51 @@ #define blockSize 512 - __global__ void kernNaiveScan(int n, int twoToPowerDMinusOne, float* odata, float* idata) - { - // get index first - int index = threadIdx.x + (blockIdx.x * blockDim.x); - if (index >= n) - { - return; - } +int* dev_gpuScanBuf; +int* dev_idata; - // then add the two numbers and put them into the global output buffer - // if (index >= twoToPowerDMinusOne) - // { - // odata[index] = idata[index - twoToPowerDMinusOne] + idata[index]; - // } - // else - // { - // odata[index] = idata[index]; - // } - } +__global__ void kernNaiveScan(int n, int twoToPowerDMinusOne, int* odata, int* idata) +{ + // get index first + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) + { + return; + } + + // then add the two numbers and put them into the global output buffer + if (index >= twoToPowerDMinusOne) + { + int one = idata[index - twoToPowerDMinusOne]; + int two = idata[index]; + int onePlusTwo = one + two; + odata[index] = onePlusTwo; + } + else + { + odata[index] = idata[index]; + } +} + +__global__ void kernShiftScan(int n, int* odata, int* idata) +{ + + // if your thread index is 0, insert a 0, otherwise everyone else do their own index - 1 in the data array + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) + { + return; + } + + if (index == 0) + { + odata[index] = 0; + } + else + { + odata[index] = idata[index - 1]; + } +} namespace StreamCompaction { namespace Naive { @@ -34,7 +60,6 @@ namespace StreamCompaction { return timer; } - /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ @@ -42,44 +67,33 @@ namespace StreamCompaction { timer().startGpuTimer(); dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); - float* dev_gpuScanBuf; - float* dev_idata; - int nNextHighestPowTwo = 1 << ilog2ceil(n); - cudaMalloc((void**)&dev_gpuScanBuf, nNextHighestPowTwo * sizeof(float)); + cudaMalloc((void**)&dev_gpuScanBuf, nNextHighestPowTwo * sizeof(int)); checkCUDAError("cudaMalloc buf failed"); - cudaMalloc((void**)&dev_idata, nNextHighestPowTwo * sizeof(float)); + cudaMalloc((void**)&dev_idata, nNextHighestPowTwo * sizeof(int)); checkCUDAError("cudaMalloc idata failed"); - cudaMemcpy((void*)dev_idata, (const void*)idata, nNextHighestPowTwo * sizeof(float), cudaMemcpyHostToDevice); + cudaMemcpy((void*)dev_idata, (const void*)idata, nNextHighestPowTwo * sizeof(int), cudaMemcpyHostToDevice); checkCUDAError("cudaMemcpy idata failed"); // call the kernel log2n number of times - bool flipped = false; - for (int i = 0; i < ilog2ceil(nNextHighestPowTwo); ++i) + for (int i = 1; i <= ilog2ceil(nNextHighestPowTwo); ++i) { // call the kernel int twoToPowerIMinusOne = 1 << (i - 1); - std::cout << ((n + blockSize - 1) / blockSize) << ", " << blockSize << std::endl; kernNaiveScan<<<((n + blockSize - 1) / blockSize) , blockSize>>>(nNextHighestPowTwo, twoToPowerIMinusOne, dev_gpuScanBuf, dev_idata); - // flip flop the buffers and keep track with a boolean (flipped = true means dev_idata has the latest data) - float* temp = dev_gpuScanBuf; + // flip flop the buffers + int* temp = dev_gpuScanBuf; dev_gpuScanBuf = dev_idata; dev_idata = temp; - flipped = !flipped; } - if (flipped) - { - cudaMemcpy(odata, dev_idata, nNextHighestPowTwo * sizeof(float), cudaMemcpyDeviceToHost); - } - else - { - cudaMemcpy(odata, dev_gpuScanBuf, nNextHighestPowTwo * sizeof(float), cudaMemcpyDeviceToHost); - } + // shift it and memcpy to out + kernShiftScan << <((n + blockSize - 1) / blockSize), blockSize >> > (nNextHighestPowTwo, dev_gpuScanBuf, dev_idata); + cudaMemcpy(odata, dev_gpuScanBuf, nNextHighestPowTwo * sizeof(float), cudaMemcpyDeviceToHost); cudaFree(dev_gpuScanBuf); cudaFree(dev_idata); diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index e3b3268..455f5e8 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -6,6 +6,7 @@ #include "common.h" #include "thrust.h" + namespace StreamCompaction { namespace Thrust { using StreamCompaction::Common::PerformanceTimer; @@ -19,9 +20,13 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startGpuTimer(); - // TODO 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::device_vector dev_idata(idata, idata + n); + thrust::device_vector dev_odata(n); + + thrust::exclusive_scan(dev_idata.begin(), dev_idata.end(), dev_odata.begin()); + + thrust::copy(dev_odata.begin(), dev_odata.end(), odata); timer().endGpuTimer(); } } From ca2fe70be379626021a2987fb8841e6698f20340 Mon Sep 17 00:00:00 2001 From: liamdugan Date: Thu, 20 Sep 2018 23:40:59 -0400 Subject: [PATCH 3/4] Modified readme and added images --- README.md | 107 +++++++++++++++++++++++++++++++-- images/BlockSize.png | Bin 0 -> 12179 bytes images/Code.PNG | Bin 0 -> 38773 bytes images/Scan.png | Bin 0 -> 14595 bytes images/StreamCompaction.png | Bin 0 -> 15625 bytes src/main.cpp | 40 ++++++------ stream_compaction/cpu.cu | 9 +-- stream_compaction/efficient.cu | 17 ++++-- stream_compaction/naive.cu | 11 +++- stream_compaction/thrust.cu | 9 +-- 10 files changed, 150 insertions(+), 43 deletions(-) create mode 100644 images/BlockSize.png create mode 100644 images/Code.PNG create mode 100644 images/Scan.png create mode 100644 images/StreamCompaction.png diff --git a/README.md b/README.md index 0e38ddb..c80fa78 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,107 @@ 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) +* Liam Dugan + * [LinkedIn](https://www.linkedin.com/in/liam-dugan-95a961135/), [personal website](http://liamdugan.com/) +* Tested on: Windows 10, Ryzen 5 1600 @ 3.20GHz 16GB, GTX 1070 16GB (Personal Computer) -### (TODO: Your README) -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +# Stream Compaction +This homework is an introduction to implementing and optimizing parallel algorithms on the GPU. To do this we were tasked with writing a CPU implementation and various GPU implementations and comparing their timing results. +## Features Completed +* CPU Scan and stream compaction implementation +* GPU Naive scan implementation +* GPU Work efficient scan and stream compaction implementation +* Wrapper for Thrust compaction +In this project I implemented the scan and stream compaction algorithms on both the CPU + +# Questions +### Roughly optimize the block sizes of each of your implementations for minimal run time on your GPU +I tested block sizes from 32 up to 1024 in power of two increments and determined that a block size of 64 was ideal for the naive implementation and a block size of 128 was ideal for the work-efficient implementations. + +For the timing results, an array of size 2^20 (~1 million) was tested while changing the block size. Time is in miliseconds. + +![](images/BlockSize.png) + +### 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). + +![](images/Scan.png) + +![](images/StreamCompaction.png) + +### Write a brief explanation of the phenomena you see here. +I believe the main performance bottleneck in my scan code is undoubtedly the memory accesses. In both my naive and work efficient scan implementations I write to global memory after every level of the tree is traversed (i.e. log(n) times). This is significantly slower than using shared memory, which would only write to global memory log(n / blocksize) times. + +Additionally, in my work efficient scan algorithm I use ping-pong buffers,but memcpy the data between buffers on every cycle of both the upsweep and the downsweep of the algorithm. Since I do not have to do this in my naive implementation, I believe that is why my work efficient scan is slower. + +![](images/Code.png) + +For the thrust scan, the reason why I believe the power of two array test runs so much slower than the non power of two code is due to some sort of internal thrust library bookkeeping. I believe once a thrust function is called, there is some sort of one-time-only process to initialize thrust specific state. Thus when we call the power-of-two length array thrust scan it has to take that extra set up time, but once we call the non power of two code, all the setup has already been completed and it can run quickly. + + +### Test Program Output + +``` + +**************** +** SCAN TESTS ** +**************** + [ 12 6 12 29 23 44 41 42 46 40 9 23 22 ... 6 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 1.47776ms (std::chrono Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356118 51356124 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 1.57856ms (std::chrono Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356052 51356083 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 5.05549ms (CUDA Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356118 51356124 ] + passed +==== naive scan, non-power-of-two ==== + elapsed time: 5.13843ms (CUDA Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356118 51356124 ] + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 9.00813ms (CUDA Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356118 51356124 ] + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 8.82483ms (CUDA Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356052 51356083 ] + passed +==== thrust scan, power-of-two ==== + elapsed time: 4.64384ms (CUDA Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356118 51356124 ] + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 0.334848ms (CUDA Measured) + [ 0 12 18 30 59 82 126 167 209 255 295 304 327 ... 51356052 51356083 ] + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 3 1 3 3 0 1 1 1 0 2 1 1 1 ... 3 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 5.27103ms (std::chrono Measured) + [ 3 1 3 3 1 1 1 2 1 1 1 2 3 ... 3 3 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 5.21855ms (std::chrono Measured) + [ 3 1 3 3 1 1 1 2 1 1 1 2 3 ... 3 3 ] + passed +==== cpu compact with scan ==== + elapsed time: 12.4281ms (std::chrono Measured) + [ 3 1 3 3 1 1 1 2 1 1 1 2 3 ... 3 3 ] + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 7.15469ms (CUDA Measured) + [ 3 1 3 3 1 1 1 2 1 1 1 2 3 ... 3 3 ] + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 7.18848ms (CUDA Measured) + [ 3 1 3 3 1 1 1 2 1 1 1 2 3 ... 3 3 ] + passed +``` \ No newline at end of file diff --git a/images/BlockSize.png b/images/BlockSize.png new file mode 100644 index 0000000000000000000000000000000000000000..4539034f416ca8f8dd745fdaa86e4fd525e4da96 GIT binary patch literal 12179 zcmbVybyQSuzwbu6QA(wxLByfK0fv&0mJR_4VQA?NK{_R+K}uR$T9l9;7^GpOySvZk z_uh5Z``&f`IQOvDEcTw+&$GYJH$UG9Q&pBF#Cwbf007}DxtD4Hfc6*sJi@sLe)EIF zm=yen=AqF;9(v~K&aW&i+Td-YO6 z!`*Ou7T=q6yoGh|6LX(ERUocJr73Mt?6R6&C=J)fz(`yyBF~JS-g?Fy>Al0cEs(I) z5`9WhZz;I#JGHeP;x0z@-qRlWd|F9iX8J+E^r%O|0Mm5GCP{ z1t8ck;NYd?|Kf7Mu|{C*4}Trhott8WA$7Z%44r(8KJGJ$0e3U$`8ReRihIsaU$G{| zxPP$}@g94*tUf!}pG?K3IS|$1*8Sv2qW6v}EVZSP`B`DtEtvN!-I>Y?h1UmKU6Ohq zEj)5C%VoUnsr1HS(&+rXZV(a&nJ)r3;&8I`PZHLl`NailB+sF7=@pa2wMRH(w))SS zZ6*&!)g2}5smr|dJxh(Lo=t^qm z)~LlR?OslfO_pA5X-(WxTDE7peJaiiIF%0P-{3_oeb81fK6uV{v zY;|Q3&d&w0!|lVK2o{~q77Noq%KKMhzm?&J93h5y`M$;zbJS0v5o^aae6-38*xGGv zjIFzNlT`{L^~9YHR$k-k7f5;;^>>qFkE|x=DD~=Wao#e={)03op6u-n z|M6qsBMQHNOKYx)%kuSs)iCYrW`$Uz($ZPx@k_^}6Sma9qGVrm&sC{N)UeQ?%N;h+ z@gpj2C2F63&E<$}G9`t}`$}S$Xp$_JCx{(5LD<+ulj0QCccnkGZQqWId%7)Wkw@ge zsa4mv$y2ral*W=MrPds4n`x-IsBqRo!gkPmdBsoE^7V(2#+%l&c#2%Pu75c5mjbOrBPqp&L z0++;XlnY&wUq@rMwmFl^ZOo+_X-(QzourfgCSo7ci(5xH4`MiL%PsC+NfwNbvoQx4 z2S8 z2R*eTjSfm&d^h-L2bOh)?RBZ~Ah-N9UPos5t*rMl2u-gtyzPiMYKqXX5pS38<%Uq8S@EY+>uo3Obhl%1S zMUm5wscq5el+29Ip$OY3&4={b3Kz010-wUV-Zh1{5w^&-)aJY@4vqBvGyL#-Oz0E` zanD;Juam8VhsFHPbolt*;xPK0~L7 zeu{+-{N}p9YtA}HTIr^b&}W!RqoUEjpNs}PU>upw~=+5>rr^ab$ zF_T9cmK$g}0Pu*>hBBn3q1=gc^KW)fV0%{8qed(MFd;7@*1MlgQx@eFWFYbD+kz&2 z%kh}Lw6t_{HZNA7vAKCLrL#=EZMX6&W(IePqC(QM`JnHw{VCYQkz)iqcfR|?9?TGTT^WnJYvXZpJHwCk3$k6Qh9P0cqSL=tBcD8I1l{)I!%PU@5gmPJFl zR}jXrY(I2B+~<`i+{Vl6GMJFcUBY4w%8aY8llISgL&PJOX=R09g1Be^A0FCcKs=Wu z`@b(A&}J9PV#)aKPaKpNhMfA@7T?~>!Uup)OakKS&a_irx9+D|jv){w?+N~^>NeR( zEFe_P1pTNsxm9rcG0$Ecmnc-|KvEO}0OC4fga}em`@Z3nl*Ryp$H0@P;1^vD%ig`e zV%2{X0>G2f4y?82COW|-_MJ~=B<5YR71t*ryzzE0Yygmyv+*A^e52`ptNzSd;IpHn z-K-T?N_V60qIVrd7R5c_WAX%`RFzx?xjxx<6C2mKn2iqLM6vtUbk^*#gb$MCva20Z zq5D+;{vJJxO2^S|+9h9U_Mp#Y#N=jzXb75~ETa8a6ghj)ph z0VNDQ>ZfG&QrQQ8-V0?FpH4d$%`x?OYJ5_4%pBA5x#hc+r`eD58L9Hz+g=Xvs1Ka{ zU26mcaF4N>xaYo!mAp3`{w3DakpbWxPvbRUhm_WSnqtvopUS z>zm2n%`@p`wUz>DL%uXYbBR`siQn$VBA{EI?C+$c3L6X#qt%8ablEweI-y@C6v1E6Z+!8+qzKAK>ey5nP@03iZN$44Q zVx4l4#V_OCHTD<=;phicY#FC+#HQJNz}lOBG|#^BsR$(#Y-cG2kCud2B{KaPyEB?c ze_Km`dOTz!1RBW?xrFhCRVpQ1OLL`MCi})H#ik_1xK$qmN9-p+2F8N;I`-*n0i_rG zbZkB@Cq`jtfSgJy5kheeCHm%Q&k2+p_=tsDWR_|VEzv98l8MYhduc-Ofe9OxQi1F>NvEaJQgFV&wa;T~(; z2Fve@8S+`hN^626+@@;@p0(w5{oS!Ku9%b)` ztJ~rLz{kpo&X-HgCfa;kQa*lqfCuNRd3Z>Esua^XQQ~mRWMBVHxN) zh0Ln+?S-)$7xRfzvt940+-*3bUEt!t_6}11)MHC}S=`dw5eKkjQu99%X%vzNZXd{VKip= zpxcattb+qL2TRIX%Dlaw(2XmH(&|i=v|@=7Ds|q|&m7w2{r&xW0Uhhk2SsE(x-Qy1xsbgNBbRgzG^ioaeE3+Vo<~pA z)YcLb5gi0D2@88&o@_VAs)3@aT1#)S@o691eMdfMdZ-z~0W(C9ZaM!Sf2gp<)x>W- z_Ibek-2+i>ZikB_OG7z!kjbmnwf~qATncn^0hst>6-FxDsRc?b$SA=%IZ{@YOdET9 z?GvTK@FDT_)wU0yytCOsN7rFeX@hFb@P+ODK+ihww-B(MN=Qy#eJ?aMHFa@$c^;vL z4Qh*cHRy4#6A}|=eNJZ+bYr5UX#+(bYrJR*W!2BKfrzwJ70Q%B0G49{u$6|1lDeR2H?;@$U_ zC&uN<@vI@_&OG9zq|j0yMs;m$Y`AQW4R@sQT0~R7;DJJWH`SijU)(In>UVq?Ngje` zm34w9D0j~m6K>uU$DAf;U+}T9t*x!BY+xSIL8{+jF$|lM{N!mq!by~H*C?8jzb>*}t%Hcd}Yt7MAwaH#dZLp@gd|S=I)}xs zfPjG3)>dspk09TJ@Ps&Qd1Q@{s^_ito!``;l=W90Q*Y_(?G><}dt+#*{D^OP_;Eu9 z-#$nz0tQB|yT?Y!Z%H*7ME;~%-o=K6ZJY-0^MibaWcZi7JhT0UPL}wv0Ts?bubQum zM*S@YeA{~a=fnE0I9)bcn1E;Q4ymSw`;zA8LeD-nar=pHA1D_PD3ssTnS*>jJ|14F z76a5fjcy;KL?b@M$HVC3C&CDKE}V$@cclXC~Q_D$tjD?`AZ zRq0Ml8kDB~(Fa9yVc7%Iui}4v7*R(TWo&aVJvway0S6&U48wm?eoMXOofYWpUQkC1 z&bbGEadL2|{A0z*$+=8kB(Q%&#kghtEVwi!#WIkOwjHu9?65{MHHhz^ zo*TjFQsW_P83&qtk151_Ir;Z+ied~=bhGLZ2xnf)Edoma{H`fvZ*HX)J=i~*hz~|{ z=-7T=81eQIu~Z@2lV-l11=U9PeFkZZ`e_3_J*!&1h=>Re=?x*jX4AX4f-~@zC+#Ur ziB#9=nXVJX$j2h@#J4?fgxM*6_q+pXw!9cXb{bQ`cs)I!xJ*OOf1!T7_>mvjj)2IhzW}h(2z3 zky&zUV%KwZ1>q~^G34My=W~E&bVt+VW|>!eTwfgT`}K*?hQSO$FgBnVtxlbF@ii^{ zD89f_nXg_l{})?{+?7Cq<#5>%QvCe8^V@yj#P9Uto3TXG%EKP;-D_kyy-W9~xs5-I zU47aC{KV#0A{aspTT%y9T1ROs65gqoE0H!YdYSqAybG>!bjWknIlxYI{`z{W0GFbh zy+Bsv?7=71y!#k%Lqy0YBN2hQYu4qjoQqe|%uR+%K|1do{mtk6ZhgRwG4Ux>pJpT{ zcXPza^>uId=Xq6~8l+|m!F-N+l=gi0=5OGw7w}@8cU;S=)PjZ(0j`Jx2li;HRfzIS z$xkvQ*&g(>7}sltFyVxW{HJM|YaaTaPw4&%pI)+_X_WGaiO1jE_=5Tag~|`7nD9wk z8Gd?Z>C=4havHfWI_vj=5S4j`c>V`l*1<#Xg1ztQ=lHf|VY1%bo>pFe7`pQTByEkR zu-yAMKJ69qvC$S@9Hq!lx~^f~haXcO!3G{4tS0Q-FhV-yDbZ5D&+Z7Kybl*TBluc8 zk2gS;Rh{>IYF{n(_EhwF>EQkGW*CEk**}V zRhjf3J;+W$7OGTRQ0YCRkk21jVIpK0n{Rl$qyU|4Xz=rbYzl;ZuKO0lcK(}m=!Zlt ztX`2wbj!Ix(4S?^6LoDjz4XHr@{&`#Z&i7$kKOObE84PrJkUj{EVJz^l`Y z9ZK(fTwf0+?E#9*jsQxLiS9v|h_E)hL3#P;?CvloT-MqYtVeC0BFMa)4C14dv@XUm zN96oj$TCpwi&3quu-xJAAS1~+wD!V34jcQk>@$%$4&mN*sBN4@pEcwcG%D3JBmdyF zI{qMEEb0-l6Z_q2&#azLC(x4j1k66DXvmjoernwUvuk^4A<9dVg6QqfpFg{<46Dp=D^$tv@j>!5MjDmimD~jUNZ-9ES|WGDQ2md9(!58-HwA*$`H^LK z)<#)7pajAPWiVJ~+RCA&{;b(J)&9BPi~qPltG=hE&tp4(S0&~}qSU||<|9(?+z;rP zlnWFvaRoY}a8$y{w=rxVSGNR}W94l^?w?H{7lZIASy)*Q4-eTDS65aJ@DQT8$WP0S zR__h?Uv#21Px|=Xw;(d;eM}X%;iLNIPWY&$;Uu|C`^PwjG-MOZ>&}&$!wI^F9 zT|zsiLQc^6Qsn!$^BkHF0*EktrLd(2oSCK3pmvt`vmgDjXN_7+q;v+P25|1yQgL#& z&^=DW0aN|SW)CowIr#bc>FcNcx_@491}L5^6#u-;;#msy8|}<7&$?x8k+x8!yyr|p zmLX_6rST8kU_Z5R+XF&nzyvfO1)IBFP8!*DU3>q8>nlDQRH%B=rdDtAYZGWZbT1q` zmL2z@fPbF2Kas^`)jKLdyJu`{Ynz*1%Zb`+tYFF`CxWqt)B7LzAwlU7%w#a6{G-0+gnA!|AMw|P1ke-0FVvESc8|nX>ddqD0z5z5QXHHm6f%jRtLiYpjgU3GpMSn zs%f{?t33cSH#b){n%Wqo7my>tSewit@LcSQD$y!4F%qtsv{~B$UnJoljfuHNNG0%j zU?np%bKKC&7KDRfhW8a44Jz~hz@-475dhE#kIhz^bcPod6$vu)wQmBzLNGu>30jx= zKnLC-Fg|*$MjSithA;~Z`V!QuqvmEyI}(z?9qU&z1@SlAivn~@^7p|<@)u} zEQ9GUoR)rn`taey$;nCE)iQ)Hs+`!LV8yZFfkRa{HCX(?4!Q_cQ5&W^cOB6#f8%ld zN1KMIyViRA;^N|FFGTE#47wRp_#mY1%8^0&~3^?3}>D=2n2b0t;|_-MeNLB_)gMl^a^~DIbk#&&tHnpaz+*z9!Yx z)$O6U@?HuOUR>XM12dGL;%qfB${rY$XT5V(cCFjHf_#>Z5N@jI$v z;t~=P#>VtCR!65gcVA*2=P8n1HEXX%3b&~6%{tmy*;#rN&HX?TDO>>G7JI3pvYvza^uH-t#8 zkt7cnuOmT?tts#i6s8#h?pH&ar&EE`^foU~*xR0FlRT%CDbrcs9t|4&jE9*GRz-}E zG&b#%{B(14DXZi9XjPCxNW?X=GDhXAoS3Daauo9ZJ?D#I_gIq_hG%`|tcD%VY0Es0 zBQ!@+N?^DO;|!jHqhskrEj(7f|wID7RnRO5@0AdAn8GDFwBLQ0Iaqxe^7 zP_i)&$=7Mh%JTJrbE!|T{`CTelb|Ptp+r@3UqaAfWZ;vzzZUuTU`)j`sZThcM7;E0 z3Hm-fVlo5)Bz3FJ37KPfd3j6akYvRZ&Rx|y&!+3`D~gLf|9-%BewZO)fP=_x1)*lC zQCsE$kCHlr$7Wgm#x@U~I5`d?mPC#TZGn;=L=2ms{bk+y_}xnd0ddKXgox;|!EjOR z-!`fXj4`13aPy~I7ouffH?NX=h0wi?1QDFib9}9 z7%dUt|1y*$yTnyK@$92c+PVwSfG5ph`C2T?t#>#DY*PRv`QYtlpW|7=&oNeL z){NhZo4xnwX^&7A*w02YAx7lTooycmmO{$TXawQEG$GLZ=YOUv*J(|fxWjAw%E3OXfTdT@gF6Go$({1Z9TNsayot%>O z@a9BBg}C`M;xmqkiJ#e)>icRR2ZV1PD0Y>$&xff5&(5|mZtou5Zp0*SSyT{3DL#w} zK6eUOnq9Yk`}Ro^XZrnrZwo|R`htA#YWZ0%SHDHI0sk8rouDq*QNM>zZ8>2d!M0eVb z8X+fMV$%IXJr{1|D`|bz$@=|H5hFZhSpJcDQFmdT!c|@(W#Oy-EEQ{a!~xEU8;f!Ol%k zwlV4753a0REN-AYd{AS7cLSWthI)74_1^@-!z#{^NM?q5qg%5-c5c*I9l#=wS3Mh# zBrBr1Ha)4<_8Z?H?xXP$rE8FR6QAEXNj7<2!HgBH{qObFzJ8$0St>=SKk)d&v)0yZ zHb~m>hjgdNDp;(bvft00%H(Vq33bK4+ji19nx0I`!3dQ*UFfm>isqYS3$BN3g-Ms< zzQ^4Pg%ASbC}pyh6gKSDHVVQ2Tw0I|l=0P1^k0bIAPoe;cibP(o56ORa7{V>5UHa0 zS~K0}$$xWRaP7+#_K@4{9L((UF!rI^=j#wc5G+b5A$mt-!l5RM4l^S{(<28k(_(6utjavh`r)t54jwUdRYubUPit zl?uj#`XyMi_rQw{pE|mnv%R*qhG0q)Y5fOi(0fzZ2lW(*Lz`c0Jda1>^zEPKye;!) zlX&8jS(2d6cfB!_&oNl_A}GdI|G_n?+PQE#IG(9KWA6FheNUgPNDh0W2pX5iHR9~jvkef(jfsUq|0hu1_UEVejp zrSwJvciR$bcATOpx{o}CqNqY|!y$VKndHmmVeL%|j4>9hdurNZ1PQWu&OKSpbwivy z+RygH3gPDG-nT>HYo3C8!*nLkwANuH&e^br+sO2q*MDz9u#;{sRgx=B`bn}ie_V*O zl@+|tZvLBUl*-E|vYwgeL!r&Xyh4PAD}5ATxG%liyr}m!`OQl|n4k#%VoO+}FGkDg zpRiEd!qOlW4^_!(aHha;FFQpLa?tjGv!CDc>iHfShVPzVlWaGuQ&sg%-$}mPoF`31 zDEhscUbe?0z9w{>(TxRh>8fU#AYc0@D^Ze3zM|6)FHd6TXGxul(kO16kixz#Bo~%V zinBM07;6|)RhLTR9b}LWdbs5MmTYbBVNwXXg&*pH{Q0fQw2{GL^SPD5zNJyXa~EVl z=Jay~SO28li#NKPZYsXJ*V*Ws_tv?@H$`&fI>?Rm5VpUH9ED);hob^G7nv%o1_^N6-GbtGQ+LVLA?foAPd83Q@ zsg;U>+u7T%!}H6uV49UmjQg*ac}^qm7f<+BrD=CT;jz#T9&p4cDCjoZ=mOR{ci=J( zFA07e1q{}H%Ax}6xWf^n*qkeapGHh&HzoAnNlzR-0z2>|0)bit`HTzj#;0DUQnG}X z_6o#l_GPjc1VDDx*G3^QxSlh(I)IR(_lY<<04KZ02l&g-|6i9H_d+Ina=$K5f5IB< zdDA}ibdKzs^zZ(<%X-fKxuI{fbO<$jDRNFiSCPE>?14z{hj}k-qV}^D z4~!AX&@CFtZLKzY?Wwf1kG9q4hbc0}uC#jt1Poi71dW2?e6lWxMSAps6D~R78$yk( ziQ=+a$&lIm%{>n3m$+Mj1sX=9h^7HZ8L{W#h~mQ9_Wn#nZF%|R^8qVSw$aT!s{Mo3 zI-5yNB)x6?4|E1bwiv56;eD(Zv^af`Q9*2xeTmr<=s!Rfx*T2X>n@L;{ccTCG2u=e zq|S-4I#+JK>6S95xt+HYN}n%m&M%gMzhK(dz09GI*cjfku3{c^5_9;q)tvkO?2NSY z8+9D}FS~%;mL?*X-}D*)w*NQNh5sKziECz6g5!Vw@b{|J z6qx+w^#u*{2=63*NS;(}Q)sC7&_j(^6ge;uZksKB>`aUP*H6$ zH27mHN=32duwC;)(Azd(D7@4f6+PSJwmUJ-Oeq?#$0N$U14BvQ78oeQIG#$R=BDRc zSwm>%TEHB^@NY+1->S&~7Fbu-<<0Q*1CYCLA-cXfMY?3GCchF~p_)^3Z=*Rw&bl1T&th zF84GBQToYgxmy8f>RmCGlXK)j9;SP7B5_Z_>7h?Y+wMVP%x80bW$*qflge;S?V-Pa z|AH=zg@pw^^_I(f$0>`0lk=UXv8A19I9bF&vsMFP3w}HbZmT34e?;XSefwN>CwFQp zf>UDbWp`s!(9U%aAEMJhYLoH9qpXrSX1LoS#*jby#!S7tc`ngG8t$m3pS-O&lR^ZPeUlh`Di2?G`81Z7C zE^BKc2yoMcz`tITLF9_-#%O1zH@>QwvCd(HIY#u#3%sr!n;Nx5>AtIU!{Kl0&=#hq z#7rw2Jd~H`rw-ssqngS_lscBFwihf>~qrdwrp7iMBNqSlAzEUWE=e+b4p!>hJ6!I?`vf6(QqjF=%)+h3>Nl%kz5kAbOn86sm0|wdN%0Qz_SZDqejR zdorosms=f5WUXIlAQEZ41XoTD-b##{G{jMUgL9@aG9GC!MX(SV3$`e|9~ytC46TSF zD>xlh(=pNRqoo&g+ycvnXBwjfSm^I$nn|7nW6DqOWEu?hA zt}l>ds<*z!W-h+vpV~Ck(zC81F}G!?z9qjQze}wfD5m@u%sE1~qSCi$p93m+%0us8 z{57=qrWi^i@}g!{-e(lcy1XXMxfwVw94CAcbd~==%J)^zxY3?h$^k0<0=1t!8269E zBV%gpo|1uHFLQH8&(T_Q+8U#lLneN6PuEzIQZ1A8j+NcrYB2Qf!OzsN(VT3c4r#D zvt^wUk*}n*I*vyT`^bHL3(m+6(!(1IawVYBRGxk7BQp_$;hD73H*ad4Vpj7RUN96f zg5y_nqNx5}h_SSkq4CaYEA4T}$BwZM?PuZqxqUd{XcC4rN>uskX@kD7n zu2O&pnEZ_!pPa0%sVOayLT+jDr&tqwr<_^?cfL_+%kO0OZ<1S$U#E19p0`HwIgf6I zu~610oyxZ%5Nd{P)lc73o>jOa_=ou;7eZg}c9-Yc8|8#;kmn0(Q&gs;7#H6LT#lCT zeY!3>E(q7OVd>~!cRoQF=}PE)e{Za-`zWml-L>rl8<+m<>?{rfh{eG{?ElNr;~~}p i|HqT%N)JBX5(@mIgqW-nl2bX1IZPqS z%*ttOSWX*ePP2{8=KJLJe*Zq7-{<}N{{H=ae{MIo+2#4X_PDO+f zNfBRl+XlQByMNgpCL$u)y#BkXE%4(#5s{3YS1y{^hPyC_h>5Be2>O~piUZm&r)gSe z?#C^@ouJ`PvZ}t8^WsO&(+sHJ9+LBjHCv*hQMEj#|-lMEEcE%0=^g7srVtRnRR#AYwjN=?q|% zh={A))}W0)^e#$^ZZw$vM?-haGJj4_2_DsBnB{o#kdkw;s*592I#e$i&nPexPrDd>eCysS5o#Qi}y~3nH8sbEJN1dw4v4bxq^WFD%6aR{vVnO)iGnj`1+*j(`O}W z;yKHr;In?D1KZyGVXR_izK~R*aLg$km$eh``-uMb?=?2!A@;x zehMKN#nXDTR<#B0;y?~fX`jX@0mdqpw<~MBJIZZwRI}94Z`S=bBre^|clewoQyUy| z^i{aP>T$t4$YQOh{f{z8mtMC*xquxz}peLEb^JU0s zf}ZJ5;+RLusup}yIufhV18+})D7^Bye(7n>p)$hn z5ogD-b~Uc;yN6!#uU>7Heesvg9;@Me*2S9qZ`MknoZ9KD!~OUr>KnuY#W8x z;S9sNY#%jA@Dl@ke5P4-b2i-c86y)aousz<^{ZppY|UwYHK>|w{OBMg{%nZ9M;-I!i z9<&rnX{onmda2tdUyuZ5Yv+WrhW-K;b-m;iet4^D;B9Ny{hOaj=pjx^ldVIiwi{aa zxyL4~Ax=zMK~j*ZRHE6h{Ay#EjZ#X(827t>Y%Xm|wp@X%*XWI$Mc%pqJmTVV%}s$q zYb61vitS^fvGIHZ(U<#l7>XR?pQ-`L758sj8^+~*$GTt9jy0eG8v}IuXK(6x@lRB) z%%mUSov*h$0=GIp26fEIXkWM$2m+?f@ZNXj(xG3|`gI3oC9_9bZ;>C&Q9wE4?W5(LF<)n?ELpB8d^SBRuhy zIPdolyNwRbaR18FTv65!JKr&+x|2O<13r>1i=SxgD0istunwobo3>c0B#v?Y1hz#1? zGFfGU8~uz|RwK&&_bxu5(cT+eNh1Wy#w?@tH7O#(#qQZEv5B>iZwB53`Q{I<_La83 z_H^}~R+Pz)d5+!ZvCD?*Ec2`k#lTT53E=!US<1!((5*p%aL-1UX|vb5t-)PRgHF0$ z7KP>54}?KPep<*&dWCc2|kPy_Goga_x?!D zEUL);&F}K`bp|beW6`cqgzjaHvkcahRy4^FK`cYN1>6>>-t&8$)LF1rwL(I!XbSAi zQiK#LN_7l=T2@t-Kp?b~87iIK&U)qFWOo}68R~po^bu_m!R_hKh`#BQF5X7xsvfuL z(~sI!vV`+%9>Q{as^*Zg&%hLIic3#uDS9Y-7yL8$Y2MIM_mIZNeK9rXI5VM>O+|K! z{z_1axkqWvxwsp5N;4}u^yCN@@1SjE1xaP8g42GTTF{Oiqc=b{^cWA&gDoN=7dlr} zM^zq2J<^MWnpSW_LS;INaxQ1&w13D6u#^rY(Rs%qc~UWJ*{Y*@{g3szV{)XwSQjfL z%3gFOVZ?^l#)t%>>e`r{$}4o$ny?ueg}f86XwC1?+F!2F1w3kxCd<-}C4AKT1NG@z z;U2kv&YT`$E>G%?Fh1M3W{a83DARecw@;Q_dHugdb72L=)1vP};D>hjs zuz&X~G{4@fk?{$2C#+H7)xzFZWq~3wHx)(yXkCFTwf2Ao~GPC@r5?Qe;F`@4OpMx$k0 zkW5p<1;wokg|y^%{L^lcXYU7ZI_az@X@fD+T6ocx=tVmunY&Ql-Re5Q z89F|=Lir{~oXT7J=2!O~El3^n7M8>12@qdwPZT-ydaK=jGu+x?ePX$BUPA}&sE5k+ zYIq|KL%lFbWn`qZc6CkrFj!PP^0H1+#N8o9l-9~uD>XT27*cLq`V_E8MILNj$exxX z{pe5~r4>yzvSYb?!a(Cih&9@GcK@ycNhs9|Zmlc(w#y^$^bt^jX@>iBO6%Tj77#xro3!%yhfTIk4y49Fv3`+3RkAcP1Fo_@5LJ5N(Yd^bf`F3_j z{fh7oiks18eL6*W$C0bsH(i#fc{;AJ@kqeLX50KzlU_>>h}$tv=N=~p)Ze&Ve$;=0 zN4JuXeL)HHi~Uro_KjTSJn01?aa5zr{kFKR4j+GhtnFnucu2B+`4g>Hrl;bidjg5B z1NG*2dB7qD2m>B`*ms#jIl~yoMT7V`MY0^-wm3tqTblB_!Bx$?C29t?fGLx6Q&TIB z2l_J}FmI;!*vCU%EJ46`d;a1yjWo68v6yU7rV?*7B+@;FvE8jds?kO_E}hdxZVCx2 z8*?2bm{C1T(Vp0|BIfUOKT_0-wg*PL zMn8}>dhn%z&yC_d<-7Y@l!reV$j-Nsfk%(IvA0FP%_j3%E^WrQv%BJa$K&kNr$3-y z&bGjZ!sf`K{b_^x$69ideM;6NkUL$*vEhlTI)}LNuN7@~$mWW8qpv0BWY7L+7z~^H z6WTA%VkkTnxCOR7DBZYzI3Tv>xm@!K5-Z1TUS{CKMo!EKC}DH7(3+6N^MAq)4GyPm zY{@cplIsyM-~{^rKB@jAjJH*2AwYWGu-Arw#RmKoXe6-KA_2z%qVa@UWsS8eLa#mS z)ya((k=~6ETSUbFqj3VOZe;x`5&2^27I2oh@PgdN+r+}i!-xl=m-hdYFKk*(KD*H_ zcX-_$fnE>A{Hx)U_^5>!p_gz8_VVA|PsfDgSHI~AW6o{dBU<>c0sOiDza7ATjOhQP zhDvK`-xJPt`C3Pmjai##cA*$L$Is-*h zC5&+k(uYFGg!yc_i7lo--|ZNE)37x5uI4*jc+DLI1mXR4?AYZt3qpnFaHdQ82$H>+ ziPRZ33OR^W?;frMAr?Mp!oYtfd{!&UsL$D%6ONetSN8xb_rx zlez5Gdm44+Z`TGQP9#jPz~0E!4HL5rxBI%7+P^ zT*H}3?G=>E`gNN2y6i5xY@?>BCd3u8ejxOYpRfPw3)%Fn=ZU+)$UmT5 zL%GF!L!sHb?~OOylitZQSZbb0G)c%lPr{TeMPGuN!}f9AyuAV-m61Vm!g7a(o1g|g z3w}#BeE3Rl@AM`3i}HJD`fIYRwXG(WNT(4}$r(WtTJ9x2; zq9S>DxN7Yl<5N+u9A}%&K7}nSP+NCj$e9}r_b8bW_*v5Le)topI1oBT+?E*4 zlbd(4cMy7N?b!C_ouDbHXwW-LZZ_P&x+h~=`MOPdYLN_^GH1C*<PO&9m+EE`IRY|@;SJ|(H^x1^HnfZMYT{>PlWLr4-MftP zDTY`ZiBRWPad+k?;x@tF#Tu1MI^Zj9$De}gS7d{#b?E0j$?R%0!?IGrp{sg{N8{w# zf_c9<^f64zw}m1Sp4;7(+&CVp&tP%hxOQguXyx~$qWy?hezJX3IG6VJT4obpGIW>> zeh6>1AkYyox}(BD8ce;AfqyV0;U4Ag{aydVtY9Rrd+2gL{+#pB34IjDyssR&tX5md zBzX;Bs>*lLJjuPVvEzfkL_-|X+?pTvmL|hD8&+m2MX3%D|BlSELai`!3P68p22v1F zA&gn>ggwk?SantKjVig2eFap(@9bc@6la{j{VLKI^dxlkSaltS=-2idQM^nUYOgL# zF3@k(jJQb*BvR>B?6l;@nt(z5vf|Kx7SlNsr@F#o+=$cF2hFI%CJ#}Te?n(QmUWc2 zKco z-XaK$*qdh?;R?vHvXD7e{e|`HsYXh1;}t-ZK>t2M5|n|27f5H~{@vhBHCbJ*$0A?@ z+4=3Wk5Us+w9$K!1NY|O1FK8ZVT4pE8INRp8fr4-f`YD>tQX z(OX0-qCYL?yG|`IL(q_~f*qp;8BA+Z2naZW;OZ<8`rp06dE+;GN=103xQiawS1^K#$A}Ixfay}W+*jAgRs$a$|%k_;O|Kb9c)k*UUDp*s6wUT}{hOMJ}C~$t5CYRfEFR0Hfa z*oaTVsxmn_Psa-y+d17Bv(-J8w1(-}Jtfw2a6IRdT57FOkIUVr+5p}&?ZLAj%&0$i zJ50ut+#UGCmv_RKr+qFU7B#r4<($PZ^KN_5NR{G;s6`bU2If@ajndr#coG-7zZd2~ z>t7nK#`!5W(#h6iuWx%pBj+4L_?eOZ<`cbaZ9z8NiVL-9OU?C$twMU7Huv|anvO}jJD?)Dj|Mk)+tW#GaG3=E_N^Jj4jcN*3(noduB&J z4;igk8Qqeg$mJwwDGMho2*k!0(#r@Ii{r(2k!{J;XVgxiv67kaLC@%YE>9*YiBUU!Hbd@qatI`VF|qM2F+&O{}FK0Sn8IKf~2 z`DEqooQ5!K)YoLtY5XNv<$Y?S!IwG`XS1Bky;boiZEC6FnLRI5N9c=Z23Kcd!lS>u zAaM?RX+FH^Nn&Ws%x*4B`KG1=U(n2MPU;NbMOHnB)i!YBtGY#P)|cG`8<}q(C7DmB z-tuXVBKC-Do}z6gyPU0%?f9+wHPoex0&ef3xSjFJ^?+6CsaGU<85vnpUyXCcRBwS9 z1*UC?Fkle3F+66PuFgHpyo$bK`F1c{r84_$fpRs zbkcArm;8dc_c3*F*-9Y{cknTO0t2m_70MSqD?7b?Dhwevj=9m3Zx12m54gC18c5CZ z6704%^FgYqkq6uimDPk*H37~2y@=)hrf0JJ*J{5lOQL!yZk5SB+9j^VzZt-qy&Zs7O=xQp z|5ej*eg%~p!97>==aNUMVa-Q!xo|UGT~k&HW}-E>67rZyx$kz6M^}wb#q3n2$S@`z!w! zxr<~I6r29)V*$s;K=Ik21o`VTGxDy0eQ5~K-|tN#4n;I_ZR-1nwUwOxG> zyH7z%AoKUQ^_}k$Hk;cXn7*?PT1%PwhpB~qtw_q9`Hn{X;huProkl}AGvu%J?=WO?$#rDtq)twQs>5BL#M5fwY zP5=)KI`%B{Ax)hs>dVT3SoGsEl>D)2NZ3`c*%T7gTJ;#i{|EqfHfW|2^O7nN$rv0OdAD0Nl zIxg3|RsHjOVRFkV9l?-Y8ka^4gH-` zm5xyu)v_%svkjE=m=W$Dr{h`gAZZlmrjc;IXRE%iPL)$_>7dPrra<=cneyyPxo2h+ z^C{>{GKn8`NT+d8Ei!|!;zMcU&==r00Wc+})_AVr)=^@esR8%uoA8xo`Kdz*UYl>q z&sw=oyA+ui`gW@^V&4|2eoz;KJ-t7<>rtS0)Q(P>R_XYQQz(0{%E6DCE=GB1b83Fi z6}I2q>*^)Mgf;`+ggJQh-9k75bbIG0=lyPUrTW?^Rh_P}o%a5=4Nc<}2AIl8jpL#4 zsl@O1-45{XX>e2$GFtyQf9-gpTmLo_yh=$NsKH@md_0B-y!csKLliZ;zTwYb0HSHG zy5wr$sJPZ?b7OVyUr_!D%D0yjl-l5hGHj;SRcL-7FwahF&6p~9)0N(C-T|hiFv;11 zGdF+5+|+m0Mqs<{$$902q18Ru9Cq6K!7p7*d?D%>`%dAm(M)Zu8w-j2BH6Zdr;reT z-^zmh1fvjBzWiEoSgxe({qT~z&4D^6RleZstQ<*DzYh}yAtf6rS_Vgd<6ZPUuQHr! z?8UW%Quej}5udi7f*<~(dG+SCxV(KHu zPFCi9$DrZf*qVKkHwC94u)j%Z-*NeVX5Gt$`SZ;-)!E`vDo&(V1M~WkhXcrSm!%nI zq3`SSL>IdDzkjH(tq+Xzq42xd-yA!|e%LwG`AR+um zufA-WuwnB3^+Rk6nM!5`MG_A8=q%p}m8A3+<)(k93?AgYq5v#Nxxk9>f?JCu_A)tD z0j1F+^%K0z^T{8q7>Vl%vypK=f6=L5OSNxmP5qMa&QoqwU}LAE8rk~y^4Wt5v7Rfo zb=%~--Jj^qF|~}&wYFRbQvlMZHq{Aoso`f^zINz<>?-uYoE&!L3_gO63cxB?=C{h0 zl~Hm+;XDkR{yo_Jq(5bF|89mHr_t6KiJm8?guwRq{S^uchTY@7sd-~dc@jCmE-^Vj z=S%-)n;9l4(b=mG`s_ulyKS66FST$0;Z5ovIR6xme6lC?XdJC_`*Cpd8?moLd`?PU z9v&B|xZOl9uRB3bCM3Gb5BJXvE~gLExo+Q;5$L@d$#);f{P_ugZ*!88#^Styp1SJ`VwX9lM<-_R&}@?dhL*qu-ne zZllip#nng0w*~PF7bj%`!5o`6rUNcN3L3)}QdMUX2S^mTIIJqM?*2x%^9qsv!}ht# zF1Va6ckzrHv(W1hvmv#Qam zA}bfQ2S&7I11U5Rbr$`Zoz1s4b49<8h8D=I?dJ9ype3=2XAl3zX~ zyv@cCpUtgk#$kHC_`Q#~!3D3B-36GcBiw_4AvI=-R?=$&U_|#S56dDF=*WKD+AvkO z5y2r5&>VaO^Rj(o{pBUBBMpsM^7_UfsnrhAPhSnxHC`>!nC!nE>Sr>zG$u>P2;lcV z9%%S96~(OIA4)l8EUQgN@?{CBwwdM%)JO0ndfsPbyQt;r#2w&c2QNt}Cs4{EYvY=2 z9(l^MX@aQ~VZ@$~NKUHWuSPD30NMuaSH=|;uojo*0eb2wb*{Pps#}@%{q|JioDkEy ztwa?qyu~7D#*w(4k})+A&;g9KDQ2J@-e>BQs`ZE(v)pM zb5CLO7gwYxt`yP%QSt+nn{@R%Y!sA?-?}DdGMFX-hqk%3PY7nkqNE28)Gki_RjWy8$@bmQmskF66 zIR~BLA7x&cqoL5^N6y~glEy3`TB>ir<+C|yF*On;Q)MsIOa@;~nCA7-&K{M59%pxy zyVTej?|J^2u^E1HKOX`{%V2uT?|~RGYxQsMZn?Gi)JM9ncETxK+vHPzF`pc6m%lxz z0{39PJ2u>Bm&>2jkV~89S2W2S;}H1%VNJ`&zcX2*(|qt`)L!Skne6*7r37~_*{|G#_+O~8ilRhL5Z7pNjTmwdIbgq#Sy~Mo1^7b7tRfI3Nv1S|@ zQ*CSyR(-djbF^WgZc3xeCuF=Mq7xXK$e)q^qX~_I(JQRqiPb6@>Xzaau$X|-PJz^D zpCW(j`X&5zHo}$nA!1e7Q=a!%OU9)Q*~CMGKB-Y9C0!ye+$+*PZmnw4%jJ&*6}+uw zBBsXhJD22vmD&UFOzw>K7*!#yXxXfxHY82}W6Uv5Lfh*N;UgJ!1{H$nnIsPfkk|&D zb-E2;0Gb?TpH|;a_iPHaz7G}-WB19^AGbRxh}F1&td*sj=YDBMIO~Wu*LKK9uj|7g z`PBzB17omLtapNJ9j|qRABM!%Jg!qdb36T5UenvQsLjw7lbikrS{g$5N;@xtGfQ~N z0qH5AgyWS4Tj%=U365^?S6}6ewl75Q=)reYDl4X~yeqGcRW@E}7HuoBBq#UTd&6!^ zmrbkwa}3k&XeL;G#y>Ft#5~&@$L+eU6Lz6Dm#tkKEcIHIrKUafYo|T7D?{cRHA*xP z>`m%Fq5q+|Ub|N*=SUV#XgkZcQl5KDm~_PjGIC|-A`j`5x(|pHrrnw}{Zc|V-tlX5 zfK$OO_M?k?BJ0XVxUi3v2AT~)_P3IEe4HX8J?NPUBP`;+6yO{v<56dsEF z&q(1vLWlpW0oY?x0M*F4Smo0Urr)?B!rd9IL`qXY>ywyKscqWTinCAtjn>Y>=gT7m z0xqG5xZC3jTDJ^`?2S%8)m$ztUzuPv`XSO!@l9A2Xh6ThhA}ec@-i}M#4I9O;a8uM zFoXq8WKZmqf1_(>2f(|Ziqk8DV863Jv~QM*rf6u9de0{JwXmfUwJ7c*E6%xupT^!U zI3aEYH$-Y-WGMBdSr3EG!x%#Vg@5*WHF^l8jEFSY>cc2EU}qRK_zyb)q84G4u1B5N zXnS@jW4a9yan8l_L@qY#b4m3KUH#lYN*a}nq-f5`bIZ&tBlD>q`3m*`=T8tN2d9Gj z{e}|8X{LSp*=N7MW3j*CNhVR9?^P!*O>SyBH=hr>?p-9lL!)w)r;5xxMGbox-e|zj zfatFAC=gvukV1y#$B)K2GUL(rLqcDj68{2~dc%n9U^~feOu3foS^8F3EUerrs_B+ssmT4><}wsX@YGdQ|8VM1Yf=KB7b}k%s$cyI@hq)s zeF^JwvKoMJ_9$DOU8d7mAgfnH2av<{t4ylC7YO5KxVMF=SL@rEeADBY&2%|b3dJ~~ z$Z_&hk|e0>0ESNLA z_@^od#LUWR7GDCJ3#7d^v$2rb9DlT429rp4c5GVP8KIF-7ZUG}a>(@SZb18Y$BTua2?O+6+7v5wo60ojt*O zG#kKg_~5KMxmY4uJIXZiS^(|BwdgLdYkI;f*ZbTk*2vnSH~Yo)52G6nPVa(-)rRn9 z&eqDjvr~+InImWRzN6Qs?R2gEM?i;j$^DEHkl{Aw=;7XHnyfZ7GSE4K?iArE2NhVi zl2J$^UQjkBtd%<9=7{9wA{D3iZN4aEG++pB&Q0Q%o*+ui#7BAYA)itc#n}CryBsHQ zlanUg`GJSiT8(fUzVojtgJX9*e4w7Kt{Ah^pim_-Oki?fJ{RY~f1h^6Ib`iwdlAE< z8X~N@#RYfgQ0%(@L4V0!vG8*zb>J0cmM|m7Gw4zRkQa+ZncT0~^z(@cKXQln!s24& zjtFRGRIs{uUhsA_=)X7 zWiESgFu^Nh@^ixu;%DnP?#$vM8Ld-tiYhA?9QfvTrVS_|0vlR~=FPJU{x@=|XJk<9{iD?w8{?z(wxIl!lJ|qk06Mw(O=k z_Gt?00M1xh^FcF5&rC1m z&}yOG?B+|WLS-3GLlFNdb#3`dUVUf$tBFaI=*3jyv5y9IS&I)cpne@$jKbG`ZpKSL zI@n9vrAW9vH4whbY3=4cqv0BBlp8v~oBXWpRN2t}RkRJB<_&ci!rjQ4mGWw$21MYSKJ4SfcK<^aBp}+tW^{ZkdY_4yYMC3ReYaOJvMbu=cQ6;Df_I zvyV2gCo=GFJ~xDjGLMg`Jv&52zBvDzlIo0U_^2LMam0fb_nygTkjk0?oG|xGWo0Cp zx5~ZCG3R(56^x{4;=WxQdYeMk(dXT`@ce>n#?P&Xn=u}yrrNG|R+Qg~?kJRfmFe1D zQkKDRyqIGQ&a#96;FnpPDK9l5_;o9$%Cm!guOoygub;V+W82Cmz6ugpE=9zI%dM~s zGWtB-3a#7Vfm>z zbe9aqN5TW-xbn2-@w?**Nij~7E}}spGrZ;CccC3)VDCU<6^|d3i`o&B{tIeq?-~Qjb{h?iYz;%NJi7TuCNd zsb`d`$$l8wI)%1UYb$s$@nOhS9xDT{GzL$9AXE3hUJ(>K({8aAD=;Zc;utFKBIT-4 z2z)pyf-HO$e16mDdU(I0?3$N;WAo!^^|O-47F{@nf3ng6`MBSv8+OuhE| zvMLi)iUfQPo}-+jgEXdhk)Pr3mKUGMV7_~N8B$2N?K|t+-o=07UwxDMU(@5>mF#$p z;bpt87?>H`eg3K?IeA3(#dZWnIF}FxR*x4M-}d=NL*4EcX4J~L!?(b4b;ibr+^qyT zu&+I^O^#DupR=KL(@(=~HD|nR!7=4)WP0@x5Co!^k+J&|#TbGCLFz(Y(TanOb>Q}# zu#rksi&JK>_6|v>sLc^W=)n&;!S5LivdrD(pO^;a*105bePOE4sL>RwSSHkM7p~=s zwn1Dt+_>93T6+$F@9KSBs`R5gJ6)>#!&0QCp#-qYXW3#>#*!6tPFDNBcs>k50PY}?k=sU)^#lNNG=Ts!5BM9V@IJ@ z3nOs@ce@$*`Uq^ZVpXbd-F3}6+(AIt$Ny9~AJ|`~Ed44yCE`3 zK8@@@=}kNsuv&%t&^~&9EbW%jr8kyB-&#Q#J~~L5)cMy5vccgLs9Jt!)o;Po8}NpB zBa(I-Rsajg{KF$Hwvie+(?*pH>bQG%Ni$|;y5b@_JwT@}>D&Us9^gf$h#!Z0}z979r2S# z$$cxvmnXFRQmXuJ``n6^nX-*Kt5^dy0{q(3T(5uLqoEe%lZ$a!dn3B2s zM^YQ|y&0 zRw8R?P;|ixM4gMMFCIKYGd|JXKAXQlLWLiE7zIXHJd?&r(0Idtc=3F7;kio91$MMy z_|~AeT`9sLr88^q#D{jbuFXXObs%$LYd`9hGy%nph={@e&(v%$hkX^1f{h$8S2i${ zs|44o=Pr$1G^)t48t~WQmuH4$6{ONIt?!`u^kSz_NXeI^lRffMsVRiL#leI5OnS{5 zMmpzMM|HeAASD_6`yZv-J+rji=}QYYSTl^UbtSv=bIak2WkF`3YRyD^iC?^l6G|t< zc&=6#^LrD&f7B#8PXKRL4uD1sc2-R+;L__(6OV~MbCYV@52W9Ffl?)`L3EcP0_kg? zwUz4ysZTvc;hB4Y;6F?JlH^#;vZ{s`6HVezY?qAQ| z4etOt+lcl|aXSb5mdBN|ez~#B+g0GXFW1@&8-n07RnDvj!soae)#%dU>$rQ6!Fai$ zBy0!FZeL|4Gtt~nXXoZ&>wG=0?<=;v6BN#rc%vbtjE9nSXgB~1{h_zWbsp&}s4lc= zj8RA;qKR1e&tA36E+QG~IBOX_=$Y(3;4)OL-xDE*8*!sd^Ut~Oil4EBafedKV~oNc zpSHysIX{S~4h?yz(foE(zYdMIR=RL=bw<-XX^RFRmNLaVxDaQ(9;Cw~pcR=Ee zy{5Dq(6sG6;^%4R4Jgqy`iXLO57SQDOYV|6tB^T&Y-Q{nQuB)-62^aDmH8?lY!)MN zl&|fAF2-;F>;~zWZ8JfQ;0~4r3ik)YN|=^r)OZg(zw|&sp%*1R;yJe?`t`ziiYPof zjc<98Y^5~ETefN#kFY56YRd-c(Vu<843MtO#$wCc**?(CA`ge|=mSHtoWJ<7-0Iyn z1A+OmTl=Vwt?RyIku!Q0SLA-4fd5pNtD2&|%|ugB9T!$wo#qP4amjjKKg&%)RVa9O z&MEYexeqpy$5G4tlJpq6!fnez@e6l{jJXFF;y$vkb}>7LFcY|EBvpp?w~7+}ZxYzb zel6!zZsiGnyrag)?gl+PSivl`Gy^njtz^CS0a{lt*LGd+)SQwOZ~2wr(hH;Eo&^aw z+C@_b-TF-H2d=zs!`b6*0gTnn=`B{rLsTu0L-drM-zyH$n@?a`UNUl?u1Q*k2hCe^V@r@I+7`qFL?MJkRfWDm3F{xjPCgzS75{J$OeRySK~r7E zX^1c6Niq@}4kR-61ssh%SMo{9VhNFDFx5Vz6tOdKM>p*DaU0$V{>w-f0C8%{0--=a z<4)j#G!$ZPvN8i~a_7)OP82VwzZwfa(P!DkGYF&j!hiR9aU;Sy%VB5STx#JeGh?e- z`*k>ZpRW~-hASms@u8Gj3^Rz-fw6aYqf@!g=Sc$*H``>n1)eTDSp$vq>Aoa4(8ds~ z8Cw^6BrL55Xz%4p9IFa&B6prW?)@Z%1x}Iq>&5&wpd3(H($zZH2^>w906?v%6yEE zJ?DZvfbs-tpt@}bJlAW7^a&{aKK4qGbLJEuMO*Aqck~HphyobJ^>dBm9dELFg|%E{ z`$q2}Wax~Ym#RtvI@?)S->&M>U|d?1%qRyCi)*6vNp!+OPx;W+AP@9p-$WA=h3{U+ zH=r7VkSdj=8z{f^?&>O=%klwqIppZ3RrAd&H_8o@lyQ~osKMi*^cut2Ysu8g6fY~QOU2d=)|Yb1E^&1%{ijiwZ&0S2z*K?0h)u_VqkL@XKI#dEVhV76Mh zRHMdZ`$Ncbfiiuts$8ySGstwEQpAumnpsh&47r%-U` zEWfewrfBbwpwR|^w-EW$SF(E1ko~|B$d*2XkO2_nvb7wh(MPc`BpQaWFyvk|+k*>Ys)Km8FIJ%U&GIGD`=l2%y%b8PrSWHB5|qW_sIL3{*}ObB0P=$ z0tbrA-0fAZD~xmiO?Uw{nyIwa>Z82psO9!BEc(rFy>@n<+IkM_-F=guwIRct=~CYo zq#2H)9zpMJq5s0g&BVVwv_oQ(*N(8oVd)lz>CoX(>K4wHOzrLVw`1JrkJg;KNTcj= zaO6b>KRV}+Jaq%i^g7X!UY&U$zVd+KyW2dU%^0d|wQva`V3_opE9RwhxywMQBZ# z63-XxP^?Gg;-LjlxhMCHj#ptH{I?82@9N|JCF8%Y5a6HC--3ti`eUK57?y>ns)y7= zO6$tImL848>I^qyw`1a2U)3Dg&=EOY~pUhR)es+)?bx~++ zNP8`Oz^F5{exhFn3i_j`v}f{LeT>7@N{}>k}F@JHL z{Oy;CSw6KjXogb17N`Wg|GIC-XvH8~SrMr~t`?GJ*3sTShmQQu+Kc}v#vrRC6bPm0 zgePlj6@s;|YVEaaevw^k-_&V+b=*qk=fj+D@c!?iKyM-eq-HTd@k&53ShR58fdJFz z623nh>bYe^}BLFK#dOg$9wItTxiqf zCENL>@ycO$A<}Duz!W=Qx^!e z^`ZLxLSo^Dh7aZrV))$S6=sWVo|frY&o~z=NVLf-FyWM|lEZEYRD)ZgLpI0a9(Wje zL~VmQ@SIfF(QWA4m)9rML{Cw`oNrK$iCO@q0nmTExMA-De`AZ<81+p5ITOS1tK`MN zR|F#Bxr~|Xi~D~WJ}DA|KGFMLe&~mkromG{@SQ8Ivr~5B`rXz;(Cy%lX^5=16Wsp> zr%b);7j<-;i@gDV2NktbT*?cE`x*hma;0LVMuY6F(sd_-Wn6Y1d~*|3;`~^-3LPJ( zZ1JC>t52~dI^3sc`L2Ep_~~JNpa?yo?Ro?5=09Bl90RSjf9N;`HMLwsSV2YD<4U5) zm>WCvVsD}>a9TwY&(EWUOH3V`~hf4t?=zrP^6t9WxV@fpE z7+Megp=ckGJioYy(8C&5J5gz1v=#eH`XFylM%}{RX%D`=q>0;*M#2Nwgg9ND1eSFz zX0Tv^0td*4j=;bj7n`WFZYUS*LB)(__S#>88moRV@yVg{J_2<~Ff!W}14E?240fDk!U?~wbR=Cb$aGGO+ zq6!;nZ!#HID2J+Xl+Qqf0Q`J)1N0Xl9if zyYfo!By1CY>*(aqr74e9-Njx1B|a_>;37_Lw@bb4)A@bAIF!85X|!%^$Te0~UU*MiyyO24e` z1aCSm`D~HAO>K?dQ*PjZN*8-jk&{q>Zp}LV$@mg@-AO=E(00;9tOwojuWD$1wb?z} z+beq<{JO&)^f^_1yWLb$rZGw5y+niM{^WIZKH3VE>pXU!!sTK0+;?zgH-`AI3eX3R zQL&untd#-dC-S=vUEUlpfAo$--5XQX-2r9et+@O4koP%c;&EFZ%o8ZBC5PUhFoB#T z74$ipmqzixa~(t7H0cjd6#u;-Nb3tibRocMU>&n#cf#)@^HE}o2iwl5{yEeX;GI|NbfoeG$hsPa76M3oU ziY+OfrOOWMYKMotgD-$%CaphX*+J{=r-;KFAt2+=({)LSh{%7|`~CaOd3RSj3~7?J z20pNU-NWWv#(#l)Gy+PqDspW%nnXU{67()@Fp5W=*S|tVL}Xg}UkayRACwX6P-KX(_}@vRZf##krNw}{HahV|8GEz1W*N|qz49ZGN>6A zkX>*p3{=g&#(F^F51P(?6=Dse>SsMVEaj11tpai!cYfJXaPGF2SKZm+Kp%>RQK-+{ zKe!Bl*py!Zs{3WZO(Kr=nCvY9o$Ft`GNJJ!k-A*G_b_MoH5$i#N+{DZSJp7Yz^%Yv zvc`Z~!z<}S&ecJyCH0Wk?$D=N6Fh9SGU5V=^tq{@5bw>M?cX9YC@K?E>0 z0rR0Q_(**pVYFTUUGGMZ)Goqdz0Vo&A^i$);*+`tY|>cZgy61}NwiH^9y-2%TDVo@ zq0e+|jp2>XZT2On^7uUzc&`pycVXG@t;*%7Imd}azuMvZvG3p$J<(}PPpC?Btw(j@ zu!8@^-FpBvwYGh~=vEOCK~X`e?yX`$suGbB5j$cORHP*$T}WusNum_N0;n`0NJIog zdJUmU5u^!72_yle1tIj1011I}2i^O5-}8LW_s*O-^PQRFj5ElXu-0N_-S>4}|NrlQ zQpHln3KbWr&iOdCX6U^J&a}SoE{EOA?8W8O_ut(ig~)}wOdjoPc|VCYFONfCu^sE_ z%AKe5Ik`2i8j!@0Yt2qcyDD2;O4Wmyhw5W`5rs~v3)Lm4wQ)*Zp%m>!RgB3O-!x&xa_^CzL-0C27peSQGWda?`iI9kD*Ni;WTv zD`$>R*xrCiAj&TxmfRQA@;ObDbLzyYEM(Hv6vb1_0_N!01)6OCzy#`9uyQ@)Qzi^< zO9(JVvBf7Hlx-_vg~R=e-C~aXnfO(&`zE+#?>Vs8N*c1$CK-AC-kiX@W0KzqW39cD zY8Oz5o3xCWEZo5KvmtTXf?E%&CQlbf;-GKQ0#M$=7T- zKlA6zo9EUX%sX02;T97R@y+Ps?mkPdb?wAdn>78zSX^ zxnd!*fGAT7v4C@KGtwe3UEc~L$Jm-!zv~bBuTm~spgwVxJ39@#Q=9cFD+}o>KOHpi zTwgThs7G^3=8`y$C`Aji_y`3|n#J~J4hFG}(me_LUfE8(lpXUn=2LF!2q|@n!B`Z? zomD?rfw>)4g;hL3$o@(BYEbuvI_*8G-2Btl0^87Xvof@6O50^JsA9X7<|z&!DpAt5 z4UtWdc< z#On^_lESKdd1ut=NU6rCk^gl@-)?q!;7_y8A-1SUgLU>E0_boQJ`)As-nnzUc;~C< z$*wz54MC;T4Ofo1>#AHvtfV`Y?1C6}aEzsWOYqA+HE)7(BHpbmJyU$7HLCiI8$mZO z!5+FJ_!LC{CbUK0gj*3w>eino`QUoF)~uy?st1X1)1Uw~&ZvUb7J9!L{naI6^o9*T zXN`SASkKitc8?OEZW_%}x{Bz8wkxG+bL<(4M?Qip#NM|r^Z?1m9-b1WcZ7Qq>O0Ir zG_~fl1RUDkF6Wq?v~{CQ9^@DoB9F83t;z$RDPNmOt*EqGRJBhk?n@0TcO4L#0Gk)# zFR}?0mt&gdsV(1Ri$mfBQ;ekf-@hwcO*tY}TaVtE1?`o1_FcR-9%ECVwWw+Y6CNV7 zkD>OCy?~8O5apyI$oUY9UZ37Mdjq$r!t{dJ0fb3@T`F?%88Y$!_aRu2X-0+%AS#L5grPTrh^`BZ1w5q1L1zSAg*C-Ygx6Q z&EhhVb6)OeI%zt?Qa)alBlIFp`POBY8@XtY2X35VetUn-2&T^6q^ll*DwR?4NLCCp zHev;MfjTZ&x_0=v6LKHl=OHrz>YvMe_lAOx0CWGlX`V{RX85V3vqm%I(P$hdAJJoHcr&@_w0x_5J8X*Tk*zcQEls2W$N+&8r@_<=FN{mFqVwZ4nx3$}ZV!oFM z^qpeq^6Xz7;a;4FsLD!H@BgOE(T{4)y@(?W|1xe~iRS5lwkfK0N=lzqnfl6xS2}99 zMN{-eT3doLoVH?y+LqZrM(OXgTn2YMsHw@SLwdNN;U<1b&Y{S zWx5y*q@?p%0Jd}9_?0wfaVp6r5pbRkaJkJ90kYxP*f$ESgL;p zy5;3}LtVBJO{}B>ypiJ~FzE_ood1UBsDNXN0n0h4!2Q_!w4|82IDz;v4AFPqmuyx9 z#d{*3(ixW7T8BbxyT{mj_LJMIv)2O8zz;*8-02j9Jke#}x89he!lD<2$IP!Xz1{jl zQ?o&H9fz3s-pX2l>Iz2di&6~GBtt};rc|VZdrK) zc3jJS3SZz+wUl4f++T1BGTRrS1H^EMiRJ8;GmfR>sio#OfO@?ZF692CNTcrz+C$Pj zWUetlu7Rk0u+dD*eS1>dEmhmtdne<<&r!k!%vm2%_7Z~9Fo!^&%#P(I zb6(4ozEwP$&p2j5e&Qz3MUL+sb))D=1gIwbfkZd@#feNKo0bZ%CcOKw@m|_B!%qaf ztU`0B1MNz|os?L(;^IL))I0@;TI_7w)II&AfAbmQAY>FX`@9AWjW-68!l=~X$9e)U z;ae{UKfcvhTf%8e{bk!b7IHUMxfU76bVtVh7-4q!YBjgt7zj-HVUPOTxo6nmQ|5wG zGsp8}+XU_3MOe32!bGDhr7f`CQR)?izc6=bW&As>rgk< zEz#@SMXw#-Zayv|7Idm3Or}5MrmSK{_HIn!+1T{-E;=V>#e`LjzFCK2Tds5r4AS$mmpoIL1n-W?nPA?S#;wDnZg>vmZXNj=RD<7?SGm z`!Hf=z|Z$FE8l3>u`&s-e5y!v=J~U)7cV6q7Tx5}cftVp zle^(4z%4}!atv%Xx>FWO*@)m7X4550dwZ~-S)S3~<&-4AQ1`QYnehArnJeEBU6Jq% zee@{kV5a5??*~JdbakRIn28i4j^RV}?4J9-+LYVgLDkW|+_ibm zqBrYpHG8jcW#%U(yM%*;$EsUj3oc&RSGvEpOI_n&h@SC%TnkmshTo+|yZrjvD&bDb z(XQdAOFb4k)RmQqQRhEo_#bZ)-AZvjz2jW0tkN%kn)kEbV1DT3+yB6REDrCCJM%zE zjlnT4ZBD0++9G8r`1(Y3(XT4ZPNp|Fq1N8Wvqvj3 z*lT+i{dkufm~lP|PF(e@Elqu^d-`gi8%+DM{??(CW~_gDu&Qw?M{8^eoe5WXcqh23 zRy=x24>JjX31b;$Qm&TbWZ?ltWYT0$z*M}Kxdmk)>draFWH#)Z;1DIuG@{h|&bf2W z`A#(*^wll!iKByw*$z(Z8oAQJj`sc|%^x=Eu4_j`Zixpd?!5XN3skF<>Wi$K{p8L3 zgv&RYYE8|+!Ag0fRuS1n%y(RcF)Xg4ikm35$b)8QG1ul*fuem`z) zeCJy-DL}r_4N;KsH>*(e@I(7b5hf#jl13vdmY7clnz}i!B}%1h8O#U%Uas3+eBT;C zljWiK%zG4=;9!{dGo1Ic%qaV6pLwxwQ&?4+V71#~7mMwFyCC)ex~VS`+`dF+;HD)P z^Fs$~tM0-%7B%w6A#^E1hB2x*(X(kYWL2jRWqc~6G-szPIybRutnSuU9q?G|hHNuV zQ~q!+A{$qe6ZX=lMlzO9{efJUMcv zwl|xuGx``aW&0rP{nVSjm@?Lj#DpTLmEl2VQ_LXyf-8Ct)1=pp`2w&t9%&s1OLs?4 zb|DP=Cq{TA!K%;wFg57$f#q6Z0pTUD&)dnQ@8Q1-|7l_W(;~|(D7T&lnOVzOX7aQQ zRWn?;YwT{u!BinB;pL0UoRf3~2_Tz4=rc37bfVW>qF@a#aaMkb{jL97DTUkLje&pU%Ky;| z_5acTXe*FiT`g_7G$4H{JSDpCW_ZO?jyEflxiHMm^X|7#$kM_tI!m_(8*(q3)~p_o$+P4JISE=kXiZGdnlm7EXrt$$;g4CN*8!QRPHDcM^IbF;<;o9 z60zTkhf`xI%c(or5F6*Le{%d*=}eY{+DsAMv5AZxaUqyhHV1N(Q9J+^UylDW^F) zafC4x>V3D!<{J`c;idjCs}NX{PASWbs<{mxkMc%8t&cqA_(#6eCK;ODzmTX9Tg0xX zK8;0EnWpT0FR&ug2=da^tCkWmICmUt%Ex-()#4EPtAb3}QGzfJq~&{4tTqbgUMPRG zD{k4TkpX_I>Q%~~A;nv0~Q zgy*z}-V9gcP0^+FtCk&gN(+PsSQq!vaxV(+h%^9|tWfTn3Bde3feVT3+D^_RAS6$P zLs9esC@E88o3x3sikaPk`#fOoP5+Sy|E6-cx};$GiH>RZv?(FDhk9?R|cw z%+Wt|pOv@EkfkOX#Hqlas!Pl(4zqSz1*jgI=Bxyni&5SJ69vbhe7OuOdXL;^I4+-N zLLc*jKT$slxJ-V%DioVFjM5hU-76*DKNqc&6&b)E_4a|=bj7WYF?~gY+D`7=7o474 z{Ewt4GLjK;`L8@cC^V)SlxvEygsPupCQ9L{tAIXRwvY8HS2HiKR*OP8x)YkvN0q-S8MJ< zhul-^XwUUx?EFmSqBWzVNEoL5q8~w7uiEm|^`kOn)jc|b|5wP;vC*urFSKR~h*W2M zso_5rSy5pjCvT28_AO-08pbmToz*cHS)0YtBC;zS`>#i3{^sQDi)jHi_<+EGYR(f# zO>b_x&UCSQC0r@Qlb#2D;p))q@~ZV4EWdu!YaCA==K%4+V!r_+!v0dXC%RwKp()K0 zW_&zkjAzU6lQm%9_)s^u?B153>KL;;q8mHG9lSRKv0a-_ry;N$uzgA$p1;k>8zjn& z*RLfn2YXOeSXl15+2Ojpo5U_g-jp`Dgi46pB3MqxOdL!bn;>r*q}f|cw7YI#oli}a z)%Km#+pK>8DP(!d;@H+$hnsau?(l0$6{su@M)4Fh*5iTyk`-ZRZ8=oWAK zZ5>T+j9wUP(Gxw36o+k|yUkGT&Yy>LTR6Mg%^?Kz^`^k z(iRud&!yQ(!WH{g2WZpd1;-fy(@Cuq5}rK@FW>0>=~YQ=q?#d}T_v0Vfst7{>9&8c(YWSq3m1dE%C0?%CMjq_A}PscvZCNrM0P<`hjh4R4a(b zXx0=ocud_(SccW=Z(ZDoTp$P9%+MzfzcJE!rJ&V*S*tHvwm++td9&iG%%(MUSS^xI zQ1|kemK<+aO8?eN6DR{c%htdcv)J+c#q+9OCx#h~7pv^=RZ~Xo*T{iaXPB}(Oxw@A zcxtKS%D7LsIUpu_x0G2r_w|iFalj)H57$uXmR+H-?IxBsRBemV>(b9C9RL{1ZI&-F z;bZL$EpESnD0zxPZYMj#7h_$pg3L}dDSwGzd(uawm1XUEBWiiJD14vqY0pE4U$D`! z-f)#F4t9mOygY(kX$+jtXANMNoMwcZkb7&7k6Cy8I^SL}_e+o|q6c_hI~U_28YMV1 z7s37`z)K{=L>2WU^?0Zj4-z`8G{S$9n8x(i*M!S$#s2weOx107aJyHpy=9IoLuJ!x z26^COzt?reWMMd9hEC!oFEL`K0M(RNA9WN3s8m#o&vV6y-qFmXQ`W3QM^zt}OmUcn z+%4zU;FW$Sqoda;BJky!^vh!V-`XRF6#!Bd-SCvBM1Wb;`P{{ti9ile1xIv5@Vb;8 z#JmY_V>=CShh~HTUHhWV(hs>IQ?LbeJN8gLW!0|vrt`uVx9F3|?RQY{$afQq7cdo? zGUTkEsRDt>(`JB2N=bC}+i!fGEL8UjH|BEJqZ8fGJY!$l%+vF;a=G4;aiYlM%sYZZ z{6q~MzqG$Gs-E^c?%dK+G@a%>!x88ZL;%r=>+w9@a2~PB+`NGu|9ITf5XfY47K|{y>{&?W;8vUQVk!Y$>z(Ea#!? z!pOr*)1%rh{Dd0#o9uftR@kr1FO&JEen?LAB=YyGPL>J{okM-iKlD=YlO z5~}+v(L*Ryu&G>SOj*olAdb_tas+cg=itdRzw|E71OKr)xZy4W!G?MytsUPBP6eps zJ}HDiN7hIZ>;Kx^wb^Ca$pY{U{uxPLT(`3V?Hp4leqbF8&iCUUXmEZvo8>3-Uc`6z zT9@ctQvUJkUjtCVAD+Kwmj70G+%Ok`0E(VSQ(SbX^SPG5Ha z_4D|G|D$#%5&2(}7d$R$v4?7Cy@c+LaC#HXJJq{{wbECEJlSdrmiZVm>P_~2JcKj( zgz_Mii9G#ew3X*tZgbpKI{^5f%OHs%qG?*~L_v$dL*f}fpCSBPW&K{Gn2_O!e35IGlX|>AfS%8^WS~H zdm3$5Tf7z*>{ggpeLGf{zC=_sYdwxhoolVY5WFdc?R_ELnu)@kS0&zEs0XQivvC+7 zdLETu-2BKZ;$)-$s_j^^?5D9-chsa0)Pd%h(MgM+>vggyqkJC^FU1f`mN8E# zEzC?sgk4tfl-TC$`DVpZ2}Q%=6Wx-+n`3=!GOBX#lLi)V8zZvt!)n39uz_i|_KZ`d zGj-YR@nB($eCyZj?5Mt?l3vTNpVDHq6_?RFeU>-!L>!<0$Z-{kktn5#I%UeYP-m0< zONoQeYhtk-(_Io#UzlYQ4JHjf<2bUb@pT@>Q}nPIVzOniE!(clxTpnVPA+@Vc&uK{ zK1&rDr{+r0YJH+%TVx}E1lA_V)Ft^e(Zs&9l34Ys7nk8QlG_9FR>#1Gct0ki-q5#YqTt&Nc3%VwcXl6nQx#xd{SDWUzQI zN)oIy+F#n;`aOR11ErJT+keC9XNSnfLQDzft4g~M?OS2%<^;^OJZYz}_v%ey?rBrf z5lozGEWcCOHfAUCJ7HgF*PGndMFaOHhC1{ljZNN%r0nmge<{}L^l6I9cCFXFMZ=Zy zDmhIN`sTUSExz}Sqsy;c>;D>ut68#XFAqbVZ4q+u8fO1|4)8fsoI*%LJ9|>d&rda@ z?~Q}CcV5Em)5JDZbFw4br+UtvF-^`8_szrKc8fOZp#5}z0y=~kA4-e-SnE_XF)Y@D zJ2l4Mg$#~fI8j~!gv4DD+6CO@;wek#l?moNx1eBHkqAgnWoWrKs&eixjxtBa0{kr> zpkG-*Pl?vQdp2?CTR;Odmc<-CfbQ>H$O>HT);y$x^`2(^-1M=!o5jk)punklB9Lw_#o6R#rwV0V z&%*;8M64Qvq6KUMTO+B>#Uj!cS?0 z$owOI)^c7ij;<-lF zt>BO9!U_{SA@#s1d;xtol)KPi7dyhV8^?4fWD{5;-i+`y062viSZ|Fb`gHj*HE(*z z$ZJ)U1Bhj@jLynJUio+lkVco1e9K&(NxlUNY%#frCYxBlj()gOHYzSXU({_-^VRBk z{zqnx&(PEwAajuTCCVmp+J-m@LVMRGoo#c!B%LE$I!%onc%n}2&qMW%@+_KTDXwa7 z{eqTeN_~9~j%o%Y?`TxF(#u1Vm4#%&3cP&UP>Jkb5;aGlDk>irhP#pODkZbCJCPj z^UHGsgR^TVObwkoP_fS&L(R^HS6hch(1vKA)HR+BEM6-yHVY<>|Msjy>}*iw!uUI4 zeUIXB6KrH@quD@ru8OOmPT7|Fh3iop?Jo4-3Z$;mDIo%Nu9+T#UOen-n!r#*-$jwd zZo*G*Y-a9<@v(%`!GA?zEYoj}ml(K!v$7V6`LPrZ(oEhh$stzt8C{SXdf&@;>^=q! zRzgynonC&rcB#e4knK5EN2oz#HUQK_hu)f~k!Ll`{O%Ab&Bw*J=7JuoBC8dKW9Sio z7Ks)73z0d7#F7^&f-G-R1y;GU&RGG1Ht(H1aFzof8k)f#k;PApC|pcaahng<4ni*w z7Yjd5cidTYKe|?X0@MhwE{;`ohw7`91*2XJkayzmHWh(Nj_Q4l3*a$?Fo?OW~KI$9ou+{!*moB_&Gn8mB3T*`QhUgQ$M|2Bd1QC+ZlokXD1#JF zreB@+LqCF5T+xyWtu~h|DwV^hrGB#Tt?9TaT3fmDad{cCYvX_urhnCOfbQ?rX>)hR zn}TO?)y}S#ISgPYK7|gB9kwIq4eCB(X(_%~J0~GWhGqjDm`=ZktoqlL~DO zKG%o?*$L%ul$P}QE1PXasgNeb|N zFZ#SmN?cv)3HcitV0e1*;LkDho7#)oY)!X>DTL?fo-uRLml3d9J=nzj7@c}5b)qEi zhoP5fSjF;@n<<-JfS=0n#@@NcgI}*mdGA_v`%R|W5iw83g!XMZYxrN$R8*tkU{-J* ziktLqg8!Qz^b#^GdXL&dokc#leBG~1Que$!<%1rpBQ;JjHqG7j+OR*yTMp|6e9`=3 zWgnZP_{DCUb|Ze2jgAgkC>)__RfgrDB^0jdVmUsBbt2~Hxe`H%T7-=F95=adez9gR z{b72g5uOv(N)=v*uK%2yIdjN~toQzTco8tm%Qn@y?d?~x+D%?JUmW{Q%sf^f@Tyct z;*HB?Tvf~a@2$-)9~KYYzN7I(fKp&GeRmF(@m_^{_bICFJuvQ_dsQZoUTkso$D8X% z3XV>%DNA*Usb+^3Rpe8_m4ApZccfSM2~fQ>@k46$GQSYws*b0<-2P zerp#{5)JG~*83uA`OC;y%j~Jr6*7DlUb*50?cnwD2TWgM`uDJjMiczjwD_JPu^iI# zI0ov$KE2Xx(KLXajxV8KeOQrjLez(?q9|*x)1~M(+Sj4KHl7-&EZSB!$~lYdXVAA& zAfmPUArfqfElZ8c=MT7S#bbPtv3qYtVN%GF!L7G$+^)NE7r!T5wR{@03J`#`2AmZP zXwgXn%+?{rE#ITfP7bMQ)dX4F`4->R7uWUM8C4PMUe_cr6_^>>aD_~@YHIM_y5EzQ zV&1!hn+G`cDv+61Kn&m6qN{J*C1hg2QwiJ22(Y}t^?iOt1Eu~8LXfElPy}xae4@1h zlQ%5N-QTNKvoOSIy94UYuG1;?r&CPla=X75>1sZiU`9=}B);gA30U#-)?P?=E!8_M z=`_&1bP<^ntw)OTeRgXpQyPCUI>IBs#g1~MT;G7>gH^Pm8DTEW+&OWPoD{!sP;Z(>aq>nXPO#j)O5~X|jQ;jd$;^For&G7gm50OR3tMOEuj)|Y zw}v8$MiDEeqtq{-%{m`6QiHxngG<@B`Ae1}I#(Q^UGdFd`ZocpW|`m{V$Kt>lZ>ZW zUu6jr$>Yz=j8fyZw!SjY8C(QGOmm2Y?)w=iEpOX2Jm7uoS>>Pakx*-myLo=5MhdH^ zAgUMBS{A@{j-_LJii%*Td48oNbVFq>{p+ffcEjb%JWoy|xRaA_MN_7u&yeyf3?en| z0I#2W{s-yfQ97E2harxuGm71VZ={bGE{lW#ux?wC?v1@N9W}m5%-fwfjN>uO7`j!Z zOb>dv1W&5=n4yz%>n*Q@Dc;3jM%LH*ejl4${e+}ci~H$l9WRF+WNnrsW2zifw=5?W z_QjinSRL){dIgzX9o=Ju+~^T?br+a4;`pj;CAbk-f+FniG zgKaVa-7r0NN6|8MbN4N8v&xfOkmx`6r-l^i&a%ISX5a1^n;DV`Fq;c6>=mTWrp-H7 zLm+y;@tF#+8o{9Qn<4aMg{(~N5CRk@Sa*YD9v_TTGY&6*_tm%NixOmuZS}Hgd@*hI=buB`C)8F1gG=`b6WO%)-z&w;cOUz(YUf|Xb)_U`XEz()No=u@ z@OAIZ^TL_x;$1oC6X8YpH1&D-;ta8*ucEspN->^b26d}{70L50d1FiK>mMQ#fN{9? z47W1*tIY8^%*+(2*YZN?(fu-)k}TWg@ec@(EPXQCXysZdm+Ie$PuNoqV6S<;rl)^| zEce{1Q{eblo)|{|1;N{>&-itm5DV14XN4cTv}ZkT?n5`{ zQ8*_TwWH=t^|T%;Y_dX>@SBTp%5KbjO7qPz#<#>OV#z2!Ytoz4$UjGAUZ&fshX4Rb zFVyT5J&qjoL2bLS?mN{D#l8cpxdjhDfNu5x7lUqiy8^eTZw`1oE7a1@!FaE z>Mu86pf+j}VGplbm%6T1Nc&U%5{$!*=_vPC{A~ z9oA<0`H!vs^(Rc5dx24kT}r1d3uCB$3}Hpwj5d0=JT_->_5Jd<0{oHE?B+KJwSn3- z#R4ZW#{5SB_D2*^20rz(V&}BhzO+T3@z6i1II((k?RGOrE*gRJ4$1RtfWQ!x)8G_` zfO@VHmsf--ivyfe8mV=)WVr*Qo-K8h9wYWt+p}qf2Y2#SX+;(XAMrJGK`xf>h_|d8Fgd}P`nGfL!(*0&9<>okO zanqxdx3^sPo?q@8C@Oi!9O<{RHkkY)w?5We&h-Cv> z7BkUZ{>4W*ub_2_P9}hu_@F#mQ$3y`)5k$TieQIF(aUi;V$nmhR`Aq`FH=HQjhp6R zSgPOhbA^Zh3PRurG77B)91$KYXK~g;a=q=M^!^XBD9Utzm+3Sn+I>mp@Eh%g#M6^7D%5-GoO@Ao zptIuSrR+krG3`?)6OS*~-Vz*Qd3tl+?TNflrXmn*a48`WdlXjF(y9=0iLRQB6faCP z0Gn#-zmlf(fCO=ePx=y5t6n|jo^&(v`7^=MkPbslVkA~q zuD^+irxLW*+X+3nQT|@y@9M4RBKoJ@7dlC85)(-4XHWSyWZ%FGookyDD-a4{;NzD zRXx%+0K#y^b$81qLlv4iRTbC&eDN(FFhj#kU@<~NU)1EbxL9-cj3sYFz_iE*N+99Pr=g6<$i4(@bAQiX1 zLF#(5zm~yxjch&A)sYPR%sA12;ymtxciOs#9K1Df!zP~Mo^m}Dg3B}e*a>!HBZ)%E z3C7)mi;ck>cX^nJERfry@>!%;_&)7stUAtc>m(im~@6gQv_yJhRJvo~C{uRx6iP$Zj#5{VANzbgvb; zumACS*$nw7lN-`K`eqb~b+lU)8pcU{uUYltI*#l3fq)ygs;~Ptz7}x#y|z{?D{U`R z>-oY|smc-_nrc0qBjny^cuOlpPXi#nolloR16%KLk;Z47Q)=mXX54YcP@&Wd=#Iv# zT3?hr2RX}Xg>&PUtw`w64qEa>R(N6Tqi`azac?|>X9#+|;`65Yv>zZJHw>H0`y0=DJ6CQPCdP+giGaYdTc2 z<+?`b-kKVaKVAE6EIR4Gl=rLQr&Q0@Ss6~E^@`10Q*4^8+{T@7H>=1g0AXG)Vr#@i$`A|Jp#JrwWW%EJ?8sWjoOB-HnM2;6$ z3;rRcQSc8cv`aQY__2?1+}_$G73eLNhe)j)cIij4`xQ$28A0T?Z9}*JL6_1uy9js? zl~SXudw;~NJSY)(HE8&Qi4UPLU#tz4?sC@ejh~%hwib*L5((0>$TeY%+C17V<*ZXTFL(U_PF9z3uPN$dCf`y?(F@?rx5qIPr(-W6nJ-VDki&awlO4zxGgs>RnXU$QSSzZaGoY5szdwcI~ zH#Zj*qT@)uQG1mWY+Ep1UiDvogqG0DC++9@$s@*l?^C4DkeF}i zh0a{lZ6v5=Xj|DuoM9;IOJg4Ku@9a-z}&4;H%`$S{Dfh5Da$=3wa{13=LAT+ zYCsKNtTn}}k5~RfT!v+``-~?3-0691LKIW@~#FM6$4ZnwR!Ycl>7(>Dx0T;IEhHE)&_OHMmfd zEYJz84}#vQE!G@{Hz|6y@?3)85UTG0CFKqcWLl?^ci^jyx|_nZpcQ~9;B2z(1}GtZ+pMO;hxZ@H%-AeDvntM0e`K7vT~q<6wbw7dsYw zvB}F#sku~}MO|t{?m+gCR?Qr&;fNvS*q0H4Er9=-&@<9FV$~)+GHnQPtD4^7<@x=k zXyDOBMZfkdwPmd9&TtLaje?3h8aF?DIRA8A=EMk^}3oN7jLyCe-%doQcdiDk*A|CD?l%upS8r2g}we(w_vZ z2REjKtE_TloiRaH?v7?gP60di0#TZWB%K}E7%HmuwBX6RzXp%(QM)vFbZeN2mYyKR z#8~?VXHfpoO<=p!);&>hy*pORa5!SK%!^o48@sBbMk(>z%LLiB*&~~V-cp-gR|O7hH+adFwLd-R_IjZ^ z2BF|t&FYr=r)fBS9uJ*Mx&CR|=2By`(1m~ZW_+0poou28u^it5Qr6iX=jaCOY9{49F*IeX>1hL5?w^6pD0c zc|&h_Lr+jgsW7zCUeNJJhuJYdL3dVde#}iFv!z4Av;Aouwp>T!oV}I0Jw*J{ z^m&=g)`&A{EXWoc4AddUyQ^#^W6F}lSV}q2USt3>$1lPYilE8V>yMU3eepgV}3#69m zlz(jbQMlrV?lps7?JGh69yAHnJKy6Q0WMdGtHxI^(f@`hy&lNgG+zQGUeOvoPa+k{ z)y_sBeRW<8NZ`E;5Q#IuahHfrzK5 z6J#5EiH3Qk`Lf8LrNXt16m+j_f9MsV$X^Z5s%DE~sqY-ygQphLK=f5uf!5m7VgxX$ zO_<5C^7R4bL`W4La?;}JMl-o>9&Y}hgaRv+H1&mt%tM683!J?%rA&Z0sfX{p7*JyV zW0KY__6_J;2IB$9%9n1$-j%J0$&-vMy&>MifgEYTV-w;l|QfU4$#vBo<4$73S zTT75hbCIo$aZiD*5M;Or2AlPJB$A~uY*}*ISH&X{p`Yu z95bUgPWsKQH z-v_Gnr{DTRQRgR9Uq)yj*fK54xJ~7FXwG!T6V=8jEE{_^L?j0+4bI)C)ztfQgQysw zL)+NIKPR)gIs;i7gRZ?2__4+ao;kUeKH4816NRgz%H^5#_MKey5rzHXJku}}va^~{ zWi}vTw&q;<1rf|TRFM$w8z8$QWAf74v~b*7TM$ZMcx|y9idtPP#ae+k0g2Ik&TIc? zYrucg$oouh_)KR4N zP_fp!`b}QnpR>-mKuTa1J7Y>kYYWvo#(1h!zUg2lAX1Dr2OZCpzkkkTF=`at1Ral0 zuHAN#I40wZBx$~deo&llC`*S#D)x#k4qg9U=f^M1er4J4~j&W zOUSs{N&g4C<67u(Opu{RK54yrWI`2yO~=c1}Bly|xyv{l`&( zqye7;rlah4oJw=%-P%DU1k+Je%taZvmr2+@~eV7wX3!qd9=z0@Yp2Wfbf&=Gr>dzyRZ!W z5Lp>Yx1df5X6O#h)rHcEl~la8vbB@%j1VtnR*p&2+6mD)^Y^hF;e;}75C zpj`m}@b$5DshU9KHs1XKk@gS?|0BzG>uKGk*NZ#04Y}b%9CXa~=~e>5ymoXoDz(!_ zPw1!=rO7(+z73Z}I@yQ~l#FlISHpZGQ!_sNscL0UZk zfENi2=6wcZC7cKB0;uCw-mxQ(Z+4ORJPa+f`NLa1q;!oXzha5xSJ`mM|AqKQ!zE0* zVEU-+bk`ik^yxT^P|?(J%ZtGl@VEIjb+0pb(X0<8 z;qf=-t>OeHc-1VnTMUhuC;0ZH7=hw`T0&e*Zrl|d0UzQ1sNB%v?cEuV%w5+y_|}7< z9%&s2X+0-WU(v!@WAPpXj*n%vo6LoW=17TXI<|!ef#Z`_S?7%XZz)$uf3Jx(=YF%L zH7I3Gxa@rfr!qy3D?`2it08p!z6>?koTpdk`{AJ#50Ywn%ZTfN;8TY;^}dUC%r=b? zba7IE)i=r8tca2102P&fDpy!RmLQz6x(#cSC5Pm<#`9i3ox?9$8{V3ozv1?AO8B)S z%5J5jSoOVJ5^|R;K%$!B99UJGBmG^zVCk3EA6$2~57C}zFaqU|xk1Q>Uz~?+hAAEg zW#~U*FQe|Mept|k%5_baLT$V++I9H)V!q(O|N5f{zi^UN#V+6mp>q_|$^4w6q^X%B zfJY}B-5rxo(jeNQP^h@)5kJ=lRs10I)Kz!7PZB5L;>nO2$Y^r#447JrH&FxGZ7L8} ze!|5Rd}FNaJMF4Im!9vjzheBirGMu@l*j5USV7S*!d6Yl`L@5&(JxA zOl|Qlv41h+cq2v@k|Qe}^To>8PysTS{I30IVK4c@Rs)W=-B#BRQiS;K){pXcK_&T7 zzUR7$R{EKqGvK3|FI-gA&p(V#O8(oU|Ld*Z0^;S1zRy9#eyl3$S<59v;B6!6HYbi$o>}gWD%UeqqldbzXQM9U975h?E_DI$R1y30^ag0`a~xKv-s^Zyw8_Lt_fhvfK0u3^SO1TDF-a3&e=e; zI`?bNL6#C2O3SV|Gia3b*gc0Q%I3SfW&N`Xb}s*Un}`&xHC&UKqgAh&F7qU~864~j zIWG1)`23xky3eyN@qMmgIr}e-TQg|vGW7(LA{>~j%Zn3Y*be=%#DoGVS9+VjC6%*pZJ;YwHp?+MXDwM>0(@8qPlkG; zrZ8v>s`5nJwBeXJ9_fwA@|6>7X=%}S`%Mq#6TAx|MYHu5H?#yE65|{Gd?edEXC8a~ z1TLx2KK3kx{?v*9tBG*R6}2~ZMiMsc!s>D`nsjg?V*8S!QXuDo}&*p0rS*?Gk@=u9|fkjwTiai z{?zz!Jra8H*svttcfXaj)ICsMYhYdT?f5Mp0i$@)X$z)s&D?gTgLylr+s*^C0yB!M z1j-MZntajJD}M|;`1_l&se7OW&#PSqIfpau&ojIN+*ikYaoXR*X{NxXE8E+A;)_08 z<^)R?YjsqY{`n^Jo#p={l^S_aLgU*4Oyl3Qe_c?137_F>cq(}0BLncN$QN(Az!wZM zM4Sbt^pD}4wVmrlnODtH2d_}yzz09m5NDmE#a8%Te96!MQ^3tuiMARq`>Zq>p4yZ!c$Sjds#&mgmr zz>4k>?%hnO0>JYde< L>gTe~DWM4f?t2lW literal 0 HcmV?d00001 diff --git a/images/Scan.png b/images/Scan.png new file mode 100644 index 0000000000000000000000000000000000000000..7df0bd59060ba98b36d5d7937d74173d8db639cd GIT binary patch literal 14595 zcmdUWWmH^Swwp0r~>} zs3;}~R{kAt4|D-#Dj*{O238e?@TdzV$Wpw%c-s>Xrx_ zU6|PdeY6Kpw~L6Lr~+H>m0jE%?CVH+*kJnB_#ldt21rPXn+^hxwrgI~*<_xbk;SEz zn~RrDor)I?%Zj>}js?@})Te0crFv^9Q1cj}L*~SGJ8D2b`SXd+KKX)!fejG6p$A=8%CnmeuG&uU=!pEjs>g!%8ph`QgwsM`BQ$d1rXU{CZekiCdG6t=-DEjXIV|A z;;C_hs|DW(4OacuZm-arKIGeo^2TlDBYtMOK-_OV$ zxl&McIlu`3CXR+8+_QR5(zLv=J*JwIAuQj;NE^o;TxFZ)1x|HIlQc_gbds@sQ_^!L2P zfhm^PF+U(SVPMr+zo@b3Yezmp$TUvlmkyw`6r%M6r;KVA)h<;e*ISAwbpP$ib~`G_p-^(EDoCUR2+mVy_G~p+!B4yXPu7 z$vkIU-Jlk;6}+D+5-~U{6(N<@!2Amxksh{SFMMEYC|boXRA~vjvjkRCpA$OD+oxX5 zw^bBw^0GFHq)c=ZA!sGSd8p=)B^0V-1FJz*6XbWS`K7@-@k}P(sqf5XOB3iL$!zwuy zJ|Wm!rJ^FOagFDO?HtvNbm1(H4rmX|?$VCkrZ&tDY3%0d%psu$X}H)`j1nW*@gyQNpcBtSOAQk9lRa#Hf z8&b<#z_Fnr5(I8VMj-?rArjc3xv@p)sk?%WVI`rGh(AEduAq{$6~!bc?h21;c2i?W za3Pfz+i)@Um=^@sN%Z3DyBV^sE}VtJL2f{E>*DtYdq@e+hGuz^A-$uV8r+arXriL3 zYO}h;c5IQf(|3mEh2rn54zkDZha~H7n=BjVbdJDk+LaPXQZ%;N`4l6Bnk0*s)>z}@ z^>!}?*mP7HPFzsC8Hd3zwiUk6wO2Xc7=Tbr@e-R%tBsry)mvs@?5r)cYI5et>szD^>`yUSoqK9d(x=(# zJ?iBf?H9nFVW=3W$s&0eu58A}mmU`5gxl zQV$fgs!*NQaGI1?UL&;j2BVz42Del?pJWp=R^B`|Kt%gk<_7rutppd+3C5<$1P8)g9)uRS1BJRZ7bmxnTa0EAXhPZld+?=H>%6^5f;)BkiSzDJ=g_a4LB5Iz z?2nQq?jXqc_YtcNKO9Tj2)Y*x`Z0#n9X z^ai5Q=j-uSP*~X5B#2=bc$iW$(}!htI)Ht7wdesvr9Vl$7U=v0Q3DH40|REB*`^lg zd)81S$-VLh0UQB-gqWfd4=mU3T8E8K@Dd6F6Fo^pUXKGz9P)?!IqKXX90ga&wPbKH#@}Z;wkBn>9UcFkr8A2MqOqr~-74vtUcS+EGTZKrP%e>D|4@>v2J>pAG zPL1Q@pz|1N$pyn=vk>bII0m7}1x?sXiD%^HRj!@95ET;k*qU7C_%6yc1Ue@n{_4Go z?{S-W+xM0>kp<%$S}wa%!#%)s$BOBSGB4t5Euc{*^jR?ct!Q+Px~yWiX0<4sO`Te?AVmJx zejcef7}&Q^BM7};?FRa3U}?mn?1sm5U4#d!e44z9;y;QFyER}-ES29*!0*jPBZ~@R zjxie(S0M=vgiurELseBzLSSATI@(w!EU6u5{xz!S3}$D>#~9y!C;z=0#XUdjd9aXUab|ap z)o?YJ?p}HQ&X%z%vwa3~CiIY0Tv-3H_c6R$d3uak?lZ4YAF2B<=_w$8@HbFsG zk1Vb-lB1Y@Si}-ZJZDZH1*MYAz$2|MoDj;7Tj36+ZS)d#_8pa|9*la zNd)Ow+w4)YEx_yHYrhdTmD&oCdcK2Lyy{bT(1E59{$R zkBSQ!F9+Hd@I$$lV8swnBUk~?b_4e`+$Y+qb+#_KBJ=b2m|_|M!0XQUN_ufm8yLu3 zz^h=dM0~}|-0-iR8 zUY&Vh^|#NSfQu=|C5IX;9)2#7aZ*s7KqkkcTe-~(JI(|5IeT1V+9{7~0#inSO`?N+ zT6?ofv)IroG?5yRYckColT3%+)K83^GfBwHconv@lEYZzeG(HT16>XK%hEw_jJMRh z8RHEvvqeE>L4QD;YVs0QEV^Z<^(K!i@JRHuqWU88$JFg22U&doRBL+OJ@a0WEuxTA zfRS{l*1LgQ0%Katt=eAgHvo~(0tA!OhzAh3I>7RJ#xZ)GbsKcOPCw6;BD zTIuqtBY{owdmuy1LUY`vC_xkFXZE6D=6kKR^oW;Oh|1(xTxiUeS**}5-MsnhQPoiX z=WS*lLKAG-12}}@>&H?xiDmIkqlZv%a7ZFV40L)FhOvVeJztcIBOnZ=V6v$GFf+>D zK)b`>lWfw%@I_Eq_N4O?WGr5WHV!nvR5@?-rPSNiD2Zx3S>i50WCr9Vz646bsf&>0 zcRW%kkxhp-McU;;b_Xjamvgzv|yd)}f*iR`_yj4Ny+{TUbH_Zu^w_}9bU)=(w z)U#hJxlEgrrdJk1ys0cc>)QLSWLqEYW!HV+oY6zvt_;O#BqqrePMk(r85-iMP=Blmi9^STMV9oPU z)tkzFyZbc6Ai{_5Suh=ghgm+^xil0JSv)9`RMObv$#`7JtHQ$aNG0Md=*o3v@>An) zjFY|dOpoN9a*?tX_BbJCvmA`Wxi1{WI(y-0z2MJL7-XsmEQxx+A;1q%IJbGd^==a- z=es=GYkiX9mHX|vc)bGbQqXtfm6KvJulzUO1HO)z)38+OYfpYqRbIgNyTo(K2i8&9&%4A2){{-xrAvF>46W@46|t6%2z44I zA}TP>lF-#ZjGMcF0Rpdeup#nIk-&$^e);C&&WHYM!p8|VX>hK~alYFpoNz~eN`F@f z5;6?Dn)D(uC{h(k*2uZwPn6VWX!JVqZw7A*M^NAzICXIGO$Zh}o*~G_lKVK!CBFDZ z`+)IBMyq)grmC}~2ul^tb-{Gl8!Vs$#z{H1Ly`y+;r&o>we9sZh(o^a)1>F%V!=bp zfyb5!GukVQc3gicTnkd_d~FAlAoGiJC`T|oEbpt|GB51mTN+^w@c!;vezJ--8@MF4 zH`*+#o85XU#weu8dnUK)&gPtl@=1^+L zehG+2N;%Y$rtk{!ZO|R>hGrak$9aY$wO>XWEDT=LUi{sFKgkvhGC?ZG$8Ba=ZEbu~ zV*xC2{oqT1gMh#y0ms<Ii_9bKxZXnw-X(?2Q5T-xq>= zz)Xpn98D}78reY3+p@B<$qM)f4MAcIJqy$~ zoVE(PK5&sI84>k2xqlsgy77v@ZRX+5Bj+uDTID zke*)To$I7LSb{a|TQwxXJ^zyT5Egk*V%J@GMXjz_j#jVxILT7q3vh5!pA#{7nu4tC z&s9oTLsW^Ul9`iZxIbN_*uuWt@v0`kZ;(H3vp-cZOAj7YoD_khN_R0?$hjltSXB#! zS|@A{f@TGgo-x>86m@28lx-!I)%vSE9p}L4Gk9&>Up4>^JFLJGiN{eLSYx}wn+#oz z{9AYa^s3v?1~TlMWCjNRWn%xID_{tX{%Qx9?lX+|SBrmz6dVxsC0Rs{VP0brkHb=^ zzJD^rUu9Nk625cyNyCJqq4-C*@&G$=a2@vhgDoz4-W|{+;3O|#zIY@R&S1p|KT$+I z+V(WcpX;AU5>KLM{aL(Lyu4SB@br85>?S*?Bp%y3Zo{&xb*CzLhp8$6j3FD!JnV#tRWPT-!hXp;dm#)e zFn`bZPh3*qS?069mS?J&3Nn(N`|@gfdbh&rP4_q&eZS)S%a71Na{E+^TIq~z;-BFS zLgH~+%z^?f+;4LiNo8*}voI@{X3Y__Fu7tVLD%@otV$t<9|oQ(_O5t-5~3Ux;sf$x zfv8&Ccl2xv&8++^kQG<=tnyf97!a-N1`FSG#~@I~O!=w=e*m4$@B$3h1w24yrUqC($sb16wHBd%G&*#>Tu>cm2(GoLs{ zt$r?rc-T2h%-@5L38q$n;hvJ@l3kr+Nq0-a7mCHHEI}W4a%wYDy-@wEX~9llWi#Bm zzLWlq4L+K}Bu;@%1$WKbJPAEMY4|(WzvBYWkx!itm#?Wi_Cw`yJg%a0HAz~hP+q|3 zCgCuy0n>9G-s_03-<^&_{cX~9ecWWI{d?1Erv@#!Y0mv$pnoq++5fxCSiR>+OkrNf z!Hp%o5M)yJ8*X!9_g@ALE4b%d@)l&!wIaV6bol#{4}E*nDDwH=Be0tMSBVvN0|iX;+O7ym1O-ZB zSr7KL@i>B=T7NCY9QW3S=R-}BBz*ID4x)$Ok@r7}z5jD2`o9U@{~Htli+zCj?U1K$ zAG-fAuRp|=tDf>&)~MU{L$y0y^~i=y_J&}Wn{{=-kWxDjAcE*hGby4;eQM?;E)sqO zPjql)vZWPrtf90&{eNNA6K^5k9hy!^Dh*qdd>D^|{{!d^i^4^SsE_t`CXlM`^6r zsv`1cOi?POb5EzvR4aI3W9HP;%*?_BjgW3uPU2nz*ElT*Jn+Z!_S7kMms*#FT0ebn zUP)fy^dMw>X_M!RuS9upvEe2=_Lic{;*4FyyTtZLm|8K9WT)#c=N47@pN31u?ID zvN+7|utT7Ke$MkGR&hY-VG+0CPd}~aq&hqSk4XwY z?E5;3?lwgfO1-H)o^4qcj$jfh)fpokTIKC7RCD@yXP>cZ)lwXeeDe4*wOr@49*@}H zZjGlqIFJn=h4JX2L(An`jjV~s%Y@|-p0z)@ZHzY@f#qJkB$O_}=R4qAI!9s-sV0a% zKDk;Uj((d#hMMWB;H$G>x{s*2M8{-VF%k|~ajM?x0CRArSiGT+z`fcfHF&cVkVk02 zfzLMl=uY~#K0#G!wvlzvSI$PdKV-Uut9BCmO}mqrA*(wbV{nxHgJ3pBim{Q4KZg)z zauQ%){XXp5Qej5pNh7YMG{lg@gF`btP8CQ$2{+!?Iu=0gI!(g_>1+?M(S18CuVB^X zK>{5}nt}-8l10Wdg`4k~v!OVc)8m}361NhwRHa#sFP$$X_TbjoHwh@t#Q+UZBOb7S5oZPb1S?L*o?Ar0wsOJ- z@8xF5YA_=_P&?jKGNPmZrQ#L$l?#y{=q^dsyU7KOKr94Z>nuQsx&%J6`M%fNBn10n zxH*Hw@08W1mf%}8CNg|~i|X=E=SVHq8Lm5w!YldeH@DtEuzvdLuR8v#O5uMGuxV>M z6v5m7c6V5&7-BooKzUv6)$Lg!HXZT@)Oylwzf4A58l_P(SkCy{Z{HS$ zy=gK*Do6#N_)Q%xw!fF&*>0h(@MQwITRS9I? z%^^SH=L|)l(y!+1ipj2cp!6Wl`O!V?80LUZTquFtP9^i^29dMbUl215l8 zv)wN{P>MM65nh!LQ`9E+QK?_ES_?#aJjE%($Gqv`~L|^ zg)BmhAV^BG0zs03f`)(yN+6wz9R6^-fowKIQiL!8C{*6>W9$(Q=f~=89;b(K4BKi{ z8DJ0=QV{xAXnrpM5`_Msu7;RU2{~d(ZO-AS1x6_Z!<05-6|TznDhf036Qa2+kd@+l zOU%&%yErtTKHegRO~e?Q6d6tcxEqJ@wivpvX(BhuSd@iel>|w3$QbxYYDFc5CH4g4 z_piH|a7!1yJ1bknefwRfoFI8Db^Aghn0HjK9%+8Vg7^i5NU9)2`fQZ{hHCsPxO%P@ z_(lf#qC$#c!{$$*2)J^!OXPr4;QI$2B?tMQkvt3Jh;z;uRRM!!TOd#U|7|GQ-PiSG zdcNfgBc?b3I67V#PG|A0}yR2v=809L^f znyISK!K_7$e_;}RDKR@fdl!I>Wc^GBpW3GxoNoN-^qq)2;1e!D#F(j;FP`v(SP~OH z9fC|VBv(mB@6s{@l_$sHcM7P$a*YR^nAVZ1Kk9V<{}qk=uc;0oV_+?EDo5vscYp3R z%(m5%({Anvbf(5Q1{5eY0a9*3B0BLD+PkafEN3B*FM5d5xp~qpZ>e%7VScTQ-TY|? zgZ|lAZFxll`VT~85Ex%$;_eT(_kpF zVQ6l~!)rk)^*4lVk;iRu?r^xuE_tO;{3@B$upsLIsoI-{n8JZ%)d%B)i?cvBnR)R^ z5$C5y#0I&J0cASnW(Kffa<+1GQ(Nn7uwZgRH zkfq)*a}f3B#Nx2UeTR7AR|npJ^qQaAc|T1s7`a2^#^vL+zfiulLW9=l&$eThD?g@n zJanhPNYxyfT_mgqSb z0+KStW)|1;{ALWzZCf+Ex>d3~O(EO|-j4N&Q4ifjYxX7p67D^R_A5@fha#RcmOFC= zg_(!;%(tSR^?~N!MbO`Wr}Y5wkQbyQ^TsZq+AY%0VC%<&#jbszNB8RBD6WJ^@5U4$~}hn z=ywMF#cH%#&560-#&Ev^!gCa6hOo?CU`Qd%`ez`E*FkRm_Svu)VzOJ-WLfw6JE?~6 z$Lvnh^%BI7TV^Zg-}yXW`gg9eC9Y@MZ(j=DbW@rbz$uzFwKJ8<70hUN2;E{w-W41*;9 zSZ@yJU^ag~zWPwZymS6RCCawdD|K^0(lwz@D4OLk{y8OrffSW*HcA%!0NGojr0-G-ESx7$q?-ZfAGQKL*sSFAsC__%v<+yl>ZNd(KEq-Y8nf97zcI`Fu9 zQ@Sz;HNRCUx=GXHudORd(giT7DXUKnU{F(k&@+h7Ozf67Z0d?xB*lBYJf3#rCaflU2e7ZI7vbw7aQu`?td^{wu zLF$i?r*wAc062bAN5(p0xF!kQjN?`?3+p0|_Rq9jxJ!vc6ATYjj|F#)dg1oQ8fm0`UPDZtPWkUXmWXN*2&tP*8Yg1PmbkHV}!WWY*ns3O634+z+Otsj` zoz46o%=)UWKPvHl)z1MSOkwyyU|bU_QGP+vVSN9C~dq z>_Pcj2&nx7h-t~shZ1`D(i;374~9&2@zozTt*_K5f0~9BSEY{X5i9vJ=fm;G_ZlY+ zp4-}WGpr0^Q?|g8e-W23Ppv#q76Lsb_~*JeRnubgKK!6$tNLsB zHxl`3P39Bsg-gLuSL6GHlGv0Uh2?UH;lWM9KeNCOaAm za&K4#b*stYCN4Q*oAo`kWOKTC6W2@*cG^jdnOEg=vG9Zj@5SOahg@BKy$!JW>h^4- zI)o5}`d>JH2|fe;DxmT&pCP+yrOcrQCYXQ;Lhk4m*8c^7$E*bgTD;&y(l3`Udpjr zo_ZH&KG7#6m|&slva&SKoo!FM&}kn5X_=w?$#Ek$>%DmQ2r^hp<(|#=%;qrTn>@Fd z-iL>6JfDq(8;y3eP)zKQ+2(xe6?IN`&~6v8NWAE)5(uc*mzZzpb&2%aA}>mEx||a4 zE$2PbOao}5XUgUoP4GajZ>+7{Pa4l0)r+&v|#=7k@EFm(rFW&%AfGeZ~hk9jWYt7*re`(f7{> z(o7eIFXg;%dHLq!WAuC!Mce2PWxiP4ACL24i8T>QLLGpt)mrI%mU$mGaFKbBkl|A& zu#J=I+OT)J-zApl32KSvwr-X2w;C>&UkRf6ESHhvR-{hxz>)bdU>wAAg;n zuak_-V(JiLab|$C;FQ^8lMqjSvFR+vq=l$+{>J1fMtf@1WJ8rBOXz0b; z$aJpvbiE!74cng>?r-jfU)XhVOxmQ9yL7a*4&(cazSe6G$gKTBx zH_2o<%!u@~$B#U)uHwFS6a_8d;oeB5+@?hsB6l!fmA|?MBpV8u@)}UV1LS}$xy+`@ zL%XuWaiYeKrY6OEQ(@$<>%E<(3#Ka~i@$ro%;I}~Y&Q2u0Gcf>lkN~&Jxf)8|E$p_ zor7Dg9rn;xvptvt0~xbqFy59ZPiK4uITmls(e1tGEeS5=?$6g0{!>7liM}f?QV?Q> zmX`XNp|haSYIK~F$ax(H51-U%kIjAb70gy^eSh7ok15PKWGBluZU1@j`v;|p-LaLP z_lLy{A7>z^^0nE#OWjpRYjMeV3klHdMh^3W$_CoP>YGnEvWdfOYVZ^kMA=M1b5Qa_ zV=7{}6tIG${9?#iwnpc*UxDjSXq;#dOH+ zM7&Pu6;QD4;XbC(+QWEuXc1p_u!O$%1v;j{Nh$Wc`FEi zG#Qmy_=b{ZQ>kh|FFsCo;OqKNHmVL{qgr+%zq3)JfTyd^J zUh!Mx6S6ei{d)#uC36}*>MV-k)-Rg7(jUIta&VwMc!?tM&?8xCGZ}a4O1EAr8I9>a_k;G15zJ4d zLinKHiGYVa7Hb^@@UmWnibH<1@Yd@0+-Joc0S6i^7_O4yEdWYwortviJ>*$HNsvxu z$xaKCz7ES(c!k_e)4^>vVb`Pc*eZkl?2*I5pt2H=#Gq{?Wn}3TR_{BZeE`eUwrx@= zXhKUJyeJd!Q!Nxk3%|l-}IEQ;(zWlT4X(5vtQv6 zb1P|EZBUbb{+w#ha2$EDTFV}C*Fb^(1 z0NL*jJ8Omng)|RArFf~tUj4{QEKd~sZ9FPH3bhAN^!xG4fYeLWAr?Y89nCaryHjBf zb@%!{WBHomgPij{H{XRK$6!pWi&+WUNpr1BB4RbMiCVqK;5XTk4<;ChPw&+8w))Xf z7UssDF+M*Sk4?k!yuGm{bP?Q4vTeDcjhx_#jX|5*Mr!F*MxUy0W0aBzGv|5G`4va?!o>_3N{v_mnB?*h88ym-Zmo z(T^7Pyt{g)(T;6cp~QOIGq(f>*ZFp*4!dd6iR)I@OUDZzCBuVwm-!SnMI(gc2i`r& zr>h-K0eTyBqav~=N4YYxo0n82oPi))!BmaqdA{RsM@D=lOAh_9wqLGrs~#u0zS^{1 zaOij3t_7P~953@c+9JJvi7}j15V!ium36(&$2@jDf8q0RMadWaOy=KHg`4{>~1C*&$UA&K7Squ-{C0pHj` z`#*Q5q=s)`z1%aMNkXRNyjd@<2_s1X%`Y%fF*MN;MR#2>Ee~)@42PDH!jxRDQ7JWE zw`@kH?6^*ZZEjD30Dne*8mq<&sKv*cs(9&wY(U*VjH z+M6mEql*RYTd1H6l1X7bx;A_`*oYi?M!f2re{L=feAcE;PxY7*GsmR+aXP_<^*X)@ zkW=vNP`|{c$YlG>IpcYkWm$;vgie=R{sU$y__|4{KmDa;^R($=44`x1!6yF+--9!N zcZP0JWS!Y<_CStT=jI|q*6VPBQZQk_1QDOY$b1Wk#kn|@hkv)_n`2l`4%+M_+}~X= zSRX}Le~?~L49|6PcVEnL<-qJ$<@TH#^2GZ=^^3QoSq6&+rf;}Xx^ME3mPTHKt*7A+ z5kj>wKt#*t#zA5m%`hYM%h}P*jIhOnH*uIxnM6(If$FG1w{gZZ{fJU1KU-gBGp^ph}Kfa zLJl|X<#nui`vCliL-zo+kX^!8ceyB6(YPkZiulZ(jJ8{2>~V58lrL8C-s7M^%{t|> zQ}c?C?{poFT#a=AHy_{l(M<^%be^Y81hhl2{DeH!cm44&6(j7W&IFS@W)AQrH(Hp* z>jYUmdtlM_2~T4BJE3P+$U_$3@hr%zr{jmSZAU!GLmiK+%2Z+2C?D?PeeXC{A2*QE z>GFIpCy{lA4My71vgD_lN4o|0R_Q}W5WB|7?Ze{56(NUg<7(;_n~y+(aj2&>6i8aQ z`apNmx^p1eCg}YL?)%dm!o}#M2C+JX)qwXRS7MRFHuquSI@xzguqG%t*zTZ*^gt-_ zI56l;4rP$J8i+FcN1n*g&Ixx%Q zS;oCWKPs_+eIhI(rjM{Ub=}|2;ZloTfu$i$Qq7H3g1l;XKAclxPL!Lj2onsL69Xa2 zrE;uPZ@O`!gpjhy(10=T(d^&xI8T4{PO-Dd+^nX3xcbfYx_W^Ia*SNSdGyLeeZ$$aMfam$SXu6AAk(>CsA4I=!2CV2bh*;2RGB z&+nbrXOpAjbD=6qD-vCiz={e>eXG)R%?w{mY7)z7RYkW(P#=F6>^4+Ewk$^o!>3<~$6F2yL^m*`Y zjtmYiIV%Z?5Mdt{3dAe-U=TlTYBA*EaSabkKyDv>ofuxoY!~kKp2icX zWh6Iu3#A39Oqi-h5}zq6<$FyU`m-feeI1Tv)^M?cb1J1dxD@K-|5jM|^W+>@JjDN# xgLa^r{_RK~nO+K;ZT<`B58u~SZWr4x$T5@tdCbxF`=DE35+bs~<$`*j{s$AFJmvrZ literal 0 HcmV?d00001 diff --git a/images/StreamCompaction.png b/images/StreamCompaction.png new file mode 100644 index 0000000000000000000000000000000000000000..3a88b814ba9303e222322111e1039b77ec4e4a97 GIT binary patch literal 15625 zcmcJ0by$_#7bcz3AgHtw(o#xDmvkc_90Vi<1f;v9Te_6)L-!#Rq&p4`g5)8EL&tom z_x|qh>dZ6q%*-D?dOp8hYp=E5cfEUm!ODu#I9PYF5D*Y>p28;Z`jh09mW(>-5urnfZfRSVw7_MpEK-s;h{?FlY9DTnERV-*se z)2fZ#6u|Pj@kgS%MP#{@RjgfAEW4aXpW^9M-PKU`X+=~Q#qqnbeNiy+AISLDJvc!^ z^|A3qd5DCBiN$xBn4X?qZ%@T_JfHOCTI>GHv$sl488fB>8Q}tT$H<>Qe`Wxt>a8xi z1V=o3gbe&U8;g(~gA90fh=>UMj*_J|0G_I1qTE725MlE1Mnpi6p~L3@o)w@Vp&=kp z$BDE8Ke7pbG*U>K!ezm?%`iKD$d$c_R?rt}9!6WCv+;~Inm`#XF=IdPO19b6y*Rn>(*8b6 zxObwr_vzO=Y+8K`93K1ka`=zUB$y5i&!jN*7Tw9ryLEc1dE**=CS@ET8jS+9l6M&K zF%qdLT@a7ST1ehHzwNo#Uho#YJSWp+TMS)a<`9UhwrzJP^jHiZv=p}?-ff5HN``mo zlvfmTuPtkys4k)pBt|{w`Kp<-~vh0vIDT*NKu+ZD8K+Gy6-c z@!(KyK1j$E%3aJAYML-U6U!$Zx=9gOCu9|}yXjx+#lS2>yV__dEC_daeF&CA)j_TH z&ZHH{PhYr%qvesQ>^IIMQ0z%;q}31)O>_;757E6y(5khmZlzN=5XC%ckf&vnNf zU(!i}!k*Yh$+b~-!^G;<*@KOuEB2Mk_y#}6o544^;D6$SlA{j1oy|f}PPq;%mXY_D(q* z`!fUzUSrLn8SI1Srta0(@qEHodOXeOIyE!WMn1Xav)4bm<)Svq^4lk$}4JAu>!cM9#uYllFv**PnAzz~jiu`~PvxUU$e5XC_o* zwe?uHT1pc64h~v}gUN{w4m+_5pPAXTks0mPTSCzlX@=GV9)GnzSH7H4sLK{1$j8d@0J-qGzK8w}5SON>;p2p8MR2m#5d>1!QS8bSPRck$EC zBb?rD;l9_%NSX+ON8db_Lp8m&Fx4W?`^soMUc5*N5bkbF2ZQRgM@9yE%mT z@bIsO7W~I*0XE$)>_8GI3T>u17TGgM0uNVTJ?7uz<^hl5d~v{9;PFZ1V&PPOQ8%@~ zVPM$gJ2M`P=Vc?>r75JI;iHDZd6!rJyS6eZ6my*oGOj#)6s_ZF{$SibvV|o(g*~-GYz~4k1D6!TcW} z?c|aS#fll>q2{c{=_U-tKAz!Iftt7!15E=)9O$+f_S#=z^@)7<$dFf)6V`q2_E3$O z#dw{$)0+@QKJvo$F<*LRx>-?C5|t49(P4G_@d;}8E(0~Sfzb1BiH~zF_r6em<>Ckn zeV$Pf%%J1a#+WP~_?{$_3L=GL*4NH>np^jgP9frgo=j6m4~<54-5&KR3a+!SlR0&5BR}8q_Y+)NiU&#jxX=K z6sNo3uTJVrjndRwK23T>p{b>&wKfi(h3xwZp-1qk1kXQtgDWbf1Zw)=#ER3L{-jtn zCBU@wOxef&rE@8D{^wR7lgIdF)Tv#>WH0X{g|_)In<8wct;2WVeZ`%8EuU_kMO6ew z@wpXV@B~CKVqXNda*_61pGca;@on{W$(^Ct52ijI?>xRkTud}rO@GFyBfjwZiO+=K z!L8TU%~*WBO%FaqOjaw#x#+!-#&!Kn&epic3{TEn-MRyv5r#3mk$BfMkUuz{QUvBER5i#uz;laJDE1BjyF6XY29T}LVz(I!J2Q0~Jq6u+m*_V@tE_@G<+E?4`9UhW zVH|Ep;0Me_ME?Y8z$r1IE_s6kZ!%GHbS+sRu7ZpktOATqTTdGuLwEVxj%S3;9^2W z0djT}CV3_KB-!&xT};+1ZSlac%=8GCxS5jCnB*>;h&5TwUAYNs-HG=)j1p@5ywTjj z^M`ldq}sxn>{)mtlEc7(gb%e_CtnLX>*pEZl0OkKn(7=HiW02EU{0f5)CzgxaaPqw z@yz!~KU?4HVvfzheINZCJeL|XUXR%`%>;*?&W$50%BZ9Z(C^1yfxiq7-chH25i3;e zlgid8R-=WRm>y({Su5bxSxE5GQJIPO0l_>n5pfgpOuc=cU}fk9y8%}D$U`L0oFEIn z>cjAtdKt8o>?HyZ8*qBE#FWW&;?g%UZH7lf)#$>T)~&1$>8P{FI1FT)6f@7>!VbR@ z@9^}9B#yRlWo+T%qKroD%l_2v&i?cC zt+VA_Gs)fkxM}1H(9-}pB>CC=HFcYNN;^(PCp=cg>sC9g^+DG1G%q*i%-LXvxT*FB z6(2;qOdPl4`*CFhXtTX3p3950s`|$%66`TcJ{T>qe4KZ+=IhZQak(qDfR>eA$_a4} zp=E@-&1@up8fG6J&=x4~rOwL2oewQA);B*?l)O}Rq~uPW2+OKfi5jlvhdk@M=@|v- zHp#aLV#F3z3B$DMdAAL(r}ueY4~(AkI_gD{xV#L0K~3mvu*op%IPY=mBZJrA%EeuM z<7pZWHb>?$htgo(Ss&Qu$?n7<&9{0X3aiiKENh_#PIQ!M6JIs*$O6(HqM)G z<}SB`@$sVt+Q_Qq1Tmi{$lFm0=KbTPmoP)goCYjShs&`K|NYpf?*I1#b&f}mL^ljS zzr&%RpjcjB9>MY^#d5=Vu3p?1EOWT=iM~Y!a`)anC_ne($KszwsELp8%krNm5k;?$ z1}J!{WJXP|l$npxYomLk1w2JzVbWZccyRXiWh@8>2gh=vIC^@88ip|Rp)Al_vZE(8 zNFsyn_4gtOM6WPjTum*u7+J*c6D6<_%=WZ|9t~yP;80R_US8fp1+ugZ0Vl-?#1En` z@eo5p;`j)Hh4h}ULN`Y90$fW#EmwU>T>qXcXUo~ynGA%VpWpb7H#OrHje$xzCj;$p z9(Kw_a$;gW7|ehu;Ch(Pjfju(-;*^wkLxymt;d2-N=)N|`ue>`=i8yji3Pi%YRU@6>Ll^P%pD62S*Xo*@o(jK`P}C*gAF;-4!}X&_YQWh2Ca2S&~EzHf6($nMV(RE3&+fgMkZd_{KobV->?7Fg8E zB)fO7xB8F&(`C&0zyKK;xnis1_wmrnFAKrojKst!diLK}!56fx1p;NToc~(9NNc9y zij|4wjyDLCX7_KK{PVF7um0D+H|(D)L!5RPo@aJHmz@0RJ=B0}e&?;cG2;X38jzlz zUPUVMZ%fv20DJ@G_3PL2$#w=X!(Z8;qi3e?#vcnKCOuzq-q9Wi4-d_0ekbdW9CuWr zzBN;bV*D1nIr9oN{KD?IvDlr}BYz+36sgz%shN@!$~|V*;IwtGp4?xf_6=4#`A-wC zO8mee2f+4hZ_XtjOcfmOD$enZ%5|Q{KE8obwO)C}ye;_^*ODF{tguY`hkk_?7IymU zemn`&L%O8XobTuP#b4Sfh^H~`P6jt5l2l=;l^d+YDEr^o=apjB4e)3ryyn$iJPMU5 zR&yGpYT3v0a=q-lj-v<_5|`8_h{&bFW$J)!%<$L&$h!6uS(*JrT4Kh}AAGXvO;*JY zFzXpB>@N?QcbAPdy1Nr>A3P(~il)VN=#WAc^1}oBf;IVp#LFZ$1CEcsO;`rx6InIJ z?zbT~sH&B5+Es13VtbjF0dayA77r*W`)WSaS6Z>2U-KtksW&1N-nh0Pzf2W<^@+!Y zxIOeAz?w$2nde3uxF4CTgKq*x6(RVsy?n#6c^D@elD~rOK!r7rzUEpiSZ(pqQR>Un z&z6Hi>(zCgjjY&`wW65Ggmn@{BbEGLSzTRi3$vq1A^qXaD8|IdX!60(XX+W&X+o)= z%8KM>nFs!CU+`#2Yj&X7wX@A5x9mSzL^_a7!>= zkhvY**Yc2ZcyS$gO@BUKS^c23J=|3jUbCO=QX2vq{Sq?2Iv*3d{$=b|{BYW8h2m~2 z-(zSMsJ<*FT9>roVW zblq1871EfTF?i3VAw{aBa-u3IZIfq-CEX$+V7iUL6f^v+=Qr>rdV(X?qA7A4`t26Y zBR5urkSMbjF%}c~n{4vKH?U($X=^}gq$UyWSCFUIFLSR)xoVM{VHFMXl+*H=bYqrV zop=g6Aa>^OvPHt=7SeH8K+ zgqpSv&zOzwO;gMwhEJU~^yB>_^+&&Hk*N)t^o(}hKHW~zql)TGUr_k)P65OxP`(U( zJL_>C#H|rRNa)dYLYW1=OMFB^!f_WHfn#Q}>?Jbhb5qkimcZ}U#|(OKLg_(d&R~Sh zUodXghRebft>2n$n4^YrQ7CA5K0Y$S`SdBP!q=^<4-XCAA}szB_+%#*DXHlye?r1P(jSSqRSP{mefC54yw%>v z5c)yY?%#5*M%d2I&Qc91VWRsOQfftj_IuL1V+NIU0|UNLi`)NpzH4^aFqX zF<2asYpX6dH#VYI-sG}%p+81zdB95kdvAz;r%pLX^?4E%n)Mj_1saYV1Gsf-@JZbg<@-ZcDbfAxW36_ zv$F4zXVYVW!2J)77rJ>yR|5j93fOk!o2db=8rXx%1f-qhH={PnZp|m13D;t-0D%TP zqfg+krsl*S8}lIOGj1JVOq%KTfA($AH@90aiRHnAR)x#zlHr#T`v7FM5S9RdY^ApL zk-as68hGG*GH+nbd(2W4`+SAQs{S3g!W8z4;(bK=@=Xqdv>;9^H(&Pdtkl<=s^x~( zbPTMoo273<<8S~oeTXGzgeQZ9vni8ne+H00W+Jody!|`r)~4xrH zBClG3(PhuEC@E`7!#>Mk`cichvoYfG&Y9UjP zw19NGH4Ku_`ZM}fm&D1f63dj=i^bdjR4XiVO4y%$5J%xg!m37S($(Aw>1p$TGRf3eeV1Z8AC zzlduH_OvsqP0NnHsaJ{oCJ%?SPSY(=8jA!3$)`I!~eiJS$<}VR~-eX@{ood(~KFIcmXt2cVI5b6QpBr^c!8v%&Gwy z85xZ{AoEBn&|34@5>Rs|F4Fh^Ej@$`MBkT2ssSy}eQyho?ik5b@pb(J$)^38L{^V1 zH$U~w<2{1&3u){b^$6J3X|J@Bj<;jVOC zI`9892G*}Mg0kHp12IFFI_F?-aDArs=yw*&2$z@>bBysj%LIXNcx!i=g_ z8lauQvHMf*4ThqvcK>RhtqTO!H*LTd@>b+Zx6=>4pU>9?`p-K3a{6vzN;44(BYe>*cBg`d%*LQ<8I zj_De!l%aGamrFPlTC9B^Tga@D{v1a+Q^b}wtMn~?S?YhU!rS!R!3tZAou92FcAT8B z#CbeaX4OzZUY?tq8^e~6_=1=0S@cNeJip^@7DZQqOQ+CQ@Bdl^%KAVKcRv8Vm4NSY z@l-@)jWOGQI=I&n1lV&Rm2Y%%l5s3lF(i>8_Wq^TyLv}0Yk>uOs@4E~7oh$lbW4$J z`$33Vs%k&n^s}p%w&CJk0F55(=2uJ_3zhk97#BH=(6QJMu?&xd;~0}{gh@On#FHNl zAmUN$qmq0Vcd4it;M3EGM5it)Yqz_x6d*tV#^jzje}&GirJuHsn+n^7Q+*y3JcLNv zbn)LH)@PJ3ZrOr(&jF6!)0=DsusRjo8|C6Ilr?GdLQ5Wkwkl(pG#)w$@6(EN4PIX7 zO7E0Hz*nA3U^3W*28Q`M8va5j!(&TB!_4O9tC6Dxqs{RH?~|r{fze&)wsT(brevzu zjxcPK4)+JZdY~ZM8~0wrIz(hz%~%jbzHh9rcUfPN7eFBcc zo*JAtFK24=k)3Xki$mC+UO!8c0NAKmuxM*vA-MtiuPP)!WCna`wP51L&R<$N6Mv+C zgR(!+`D@WCl&!OOnoc1h<2+Ip!X|t@G}h|)b0P(Jwsjgi+E0G1VF?a%9t_z?YiS#^ zG=4+ zY-)#o-jrW(KT>)BXzcCEIIJgJmjI*hxuv?r9Ke*g=V0=F>ezDIBQQ>F}y?#R|374S0J zo$wkPhkfrCs1S9*Q2z^o|Dgq}d^=Oz55D)b#LMCP)(nXn7*w13#}`q#TpUey17#rZ z=1;QTh0kPI2x{Dc>&$sFCH!dy-fpp6=#&rO4K>Th2_@RcP|HaW^C)hy;idzh%%5Pd zbumE(vg3FO)35}cuU2?|)UVCUeP!j+^(Ro`XIKSV z#~=?Q$@l_jNP2fcV?~W1an2ubGv-#`%gyQot^-X@cyi0@g<&IYZ4Js0jX&(cqX9PK z&FahJwUQ_FV=(yHyWn6` zzXC$VNmlu{A&hFy9RvhMU8uCc_p|k{4_CE+)LGxHTHd?rU3_H-`J0>l?HNXz=MKkb zea(3GaDMqsxU5N4UY_apboI)Nr=-g718AO)-_$a4Qwwf@la_`a>EFEgyTchBojF)| znkFnE*w}$cerdUr815PFWXkKIlCNdI@aiSe3|OB2e#j%jznwKo8PVwM zfqzjgn@pZ7%wl?aT^J+L3LteRX+YV}soVTAA38Rso>T(@u;-Rl-dO0sZ(a{Q!sGTlg0Xw7_n-c!yu z5}`>PsNaSyVYAvxCfUT~r|c@z3RKdx3bsQ6Z_L+5K@tCH8;rDio5q0mMs|#tTGl!! zb^NHGB%Z!@tK+Anx-3!v4WMez{mTclxe>Es+-t-MPl5tT32y`fW$SpG4gAe&^K&fo zr6^Y(Ztf`W?S&RlGVN&1#7tQgnAs4j{u4w9yH!teB z`P!Q=?T+=amuKBDf0dd*oe>~Mtr7ApS*witXA@l5rRyu`j4Gr?n9trum5k(p-UEaC z_wQGbhgbY<%I^CHkC1*!&yAM=#g^Ks~P-VDdxW?CGC6T1rg4GKhhirx5pHC1MSdg015j)GI0E<CH5!zK9U;mB0|4$CuE^F>~71YwjV6qY4r+&5uGBR`NDPCJh9n$Y(`I=?%lPPW9HE z0@W#jYbAoE-!=&p8pdeAV|C8^hAcB+wY4vtgPJmspl@swCE{Fj*Cx3zPsgoOJxvyj zr@heXf(NC&*n@{iCb2jHlYHCjdGu)H5NtF{+I-mfV4Pju@*eEHUIs#KS>WMxDy1{} z_u0oESb9slsNIOmu*XSu-VFaZ7wYJFVUwdK5p&t_$Tp#>(CAHAmBoIFO|0gfp6rNv zFwR)b5!7n>to^%U)@J!Dy$XSFd3qv`^PDl*#ivLu5lE0f+wEikqNzoeOxiEx7*rF|cS6)Lf9=?%cO#7!?c27(6{`~ z{rU>sxVl$tq{3CkrEAVAcD3f8MCw=*W!{{y$u-7grHs{+^?|8IGF{*99kC2tY7wc^ z5jYex?1mjiAT=@~{!)qIfU>DR5B^#do2~Ic9sV>yrU+0tat^;dN!N)jD6!FleC_OK zRf^(Fn>|A}VIl_9J+8})t-4kr zJj@3se6oCUe&n>A6zexH4+y&LXk%jxXQ{Kj7`INX>Vx!-2g9yZr?I_P z9O9a^^^fU2HCY$K1(WoHQPRjx7a(EE;p&5%FW_tXt<1BwC-2Gg**@Hxl1RGa$`v3G z^avEGOYkEy=;~{=Q2?!+ae{r&k7Py(T61*o42TIiyco^#@oOMXc5{X4N*W-`p!OnF z8k951Mz(3UxTed(QXo&@ewxv)`S2Fyvj0i`En=EDe2)E>h+d7Gd)PL&BK5tLbgXEG zYqbj|*T%y4hZF9mG`F0ElTOq|=wrxqCfLls#M30FG@l<87ppF7b8l6mh-B}L(qy%C zj;JbhI-rWy)R4*%!+60y?R+y}gDZjNyf(EIqgW4{!Nq5!h$1}flcPBH!M{&T^NO@gS_`GA*w%^%ry?r8X3g6HcJp9RPU7!+%D&#>p{XANp-|D$t`YbtY3 zhk`tr&SgtQ7S97KS*iDN)=H9>Wy4rTD^{AUd5ZeU_)Asg*C9iBq_wG8Ua+wm%14Iu z-|*cErlBw`n_|oPGo^kTo<7iN{P1e|E4P&>O@m@+YpV$6&3L#?Q$BB+(xP^5>Rdi$IWZ+R2rfY^6|yp_jz{l z=q(qT9mRpxYtyA~g5txoM-Yh3Q@%-f`(j(eJ%-!m!XcTE_qzug=z$HZ`z=Q#=xMtG zRSFZ@KF#_+b)63DmKcCx#5vs#A_$49z$O#bq5E69%+^>2DhxNJIeaP&4yUOPvfN7_ zp3YX39st6@XRX*1Im%vZuKlj{SQuKUtFgM^Cu!-`b_W{8CA{o)j8ixLWqhUMVH!%Q zy!(~51B>yIxYpN)E0-XH$kSl-1n_GECXLCw!H2kbJMM?oly?I>3?+B(jaXp2;+p{a z5jY|OSDVcnrH~Z3VEa{FrWGaNnt}|IWJQ#)B*4n_5QPD^C%CAQJ!X;ao4HhqHP=;F z{`sZqfZ<`S^LKaPhMdfAWKPihRAskaqL^1`vkr^vIb8z2c@l zBmb@up=4*$-CWg#iAQnCg3PA3iE36QRuX_?_^>bxIsKS|zZlqqNbh>Dgza+yOStI- z>Xm{+w^x_@_mTGaEUky;XTs8Kn0S7Zwkg32yW`wF!p=ju2j-C0tQvEz?Xz%pUT)r6 zp=V1EglK{zyzvf+m>evG)C#bBlVs;2&ge7(j4Tz^f9c|kt90dm7d~RBVciM%Y^ux$ zpG7ujBPsX2f`SqXu(_m#yey$D85`Gc+h5bAN6k-qy#he_{Q49W7brU7kb?PgP5qw&db~K| zET)c6zR$f&^JsamBe48rk6Nu@**-}&oW+DG?j&-&oQDAmI0?gA@T7b3H-$Dv>@9g6 zb3fxT5oi#(g2}q@JBN&eO=V_4;m69F@+{8-!)Y0c`76_^(1yx`jDnnxnWvYkDZ+#5 zJtrr@V}_QB#GcV_qrGwP@UV?A9^1^%ANWZW{YC7984X-kETuf5ezwZI>NY+>S+~5T zbdmx^ne0Y-`l?)DUpe-@Ev(PZ)TTdj=7xv%9??B?OC7>Y|Mie z;N&$Yhi;ctE5_B~nsMXr-J!*N-VEG_z!fwz=^I+MUfZ3WKhJ1?G;v;CgnpGA!R(hf9Lu5>Vr}=Ay_b!rHx#Ha; zLq?nmqaT?KeFD!U?@%nUhp*$~b?Gg*EE}%w?MaTrWKyDapW{&YGpnBmZq`s1TWx!} zuZU)OU1HoWgiw9&dm;EpR65$LPZyAY)I~POCnse_=*-N_$f&3i9E3)4^ifc5`Nx7< z4YOH5W-iAK@+)tn?qX_ZikoW;BIm}eJmueQY`VsFTX3!LJXo53S7Hz15)QBt9R(=v z+GBWHT5YlPW&dt%{5EV0w!S_CZECEv9-WxZ)r$nb5pL}AN))SJ#&*K>bX@IP`*tzj z=%)c(i@{n?kNI_%X^qZdQ&OZ=3 zLFvNf@sAf?2QVy+bGbmSNIXZc-GL3Na@&r)Bf4d9)&+3h(RK!h-6*%!y`mH5*+7kn zL55)(kNNzoyk3;5bZ3|0Lx;Ro>@uQD7p;D=)X5YjCh2lIpNu8|cPckypl)V;_4aMm$B#Z?Feq5qel@I4d2jxkvC=kFweZItAtQ$Rcic2D zx(#4f_A?I;j}%@jLOVJ$E355|jXQVlV7e>M?~)ckeMl2s!?V?Z)Z&N#&s#ka?Pokf zLh1b`8_#=*_SQ&QX=&+bo|2Nhe3xU#hl>iS-)CWQs?oJfrh@M|l?@<@D8@f; zEJ@Q{GHu(28P+@6Z*I`}*?TqZ7Bu)RC5a0FRVxpMC~zz4?^Ee~Zx8`hCnJ_BxPGwd zbix-*?t0de;kHV|3&`0&!wR?>`1Agb3Vn;AJ9!qP{)xlmjR`;}78~h^6n5X2mNV8n zEB{x=K~S)rV1x?BF8QEIqa%E*h4-W{!)&u7B5TlHkylpu-?jJfCDZXT8NY0PLqgN! z!6*anxRBk^l+vahiE~%zXym__Xzm$z4lFTXb1(8Lb^2|~#f5pI52@R;EMiISkAZea?P&j)UKDrR&`M%W|*LM=xBu4Tn?c-CU>aEkyWExy+)IMSU+n?MJqw1(czXmY;q^9L0QlJDC^ym4e6BtCemN?P-&HEd2g|?=95ZA zX{~p55dX@dM9%Axac?IIegvqcsv9i92l%?@BlM})M>4qU8L|}mnqX~)IN4=*wIy${ z&h&n0-f_!Ge-(xX8?}MKZB3^bU&Jcs zN$9+FwQ@aoeZYVexutYH*e!W2F3ihvowGIk(RO@<8|!BH?#@^SH&~S^dVX#cI_7%v zi<0L~xC7K9U)XI2V0KLl;&2+iCWI3N%Hy$n^0+&*NiDda3U!|Z^yd;H)p)|~u7y3f zTV-3YYRXM@uZy~yPPSW(I+dBOQUZydjz9EzvWrC>DP|%8q zb)REl6oZGS^i~n#IOaiOMmdq5z^zx<(S*V|*ogvELhvY2-T0*I6RYR&Zvje;9K8Agek+nDQ0odVk92j4ZxckEmv6ZA9l&D7M;ks!+crcP>vCLy-2D~I1OmI z9xmOF90bdc+JRdh75u~uXViXET2#FX#X@JrI_`Hsl<+XrOkn(#Y&&#N-gvC3_H5yEfTmi~_ zx`Qj3##m(69+_GQPv(DgiXE^A^RGg|2x>|3mU89dUv7CzEO}{%3!W%FI=nsUfN}HP z9;5F)BBz#e&LX)nqoPRC1K#I2G!@*v(CJvb^esgP^{?IQs<-4Z4Z4^p=CTJ~d1cz` z)4ln!f4sf>g^K~biZ2 zE~Up0oPUTAcl3f0)5u)rAz^!Rf&5yR^`v`PPUJ##=!MwpQoJfA@8;#X@g*bY+9+oC zKBq~&y!ou~gBfsW1f6wCiil0zOK3z3Zv*WH_^}Q8Hcm4YYgydlqd|9tkj}bBRQ%O zAZE43yzCI`yOxoi#yAplUy{YzrP>};e(+xWWTBAp+)P)`_4ZID{*HgWnz6m7h3PYR z&viv0D@u}GU6zF}TM;>8EM8(EIGp`1-_;Cj^xo|XUHr=m<_$_jt7>aSY3}c z7OU+>y~Os}{CE7HKv^@burfIf-Zcj(TU~_f$%(yipr_)c0Q*VF?J}Y3wz7PnMf&nt z=|abQCA3PsHYD^-(8V&tLlTaS(8u@D=o8AqVo~}6hSrz|Hm;elUPat`K9s#E-~X6- zK}=b0im{7gwOYTF&Eh)^xasSMJQ>vaNyyyAZ4T!ADDQF+5~`5bqVDwtzG*SyJNbe4X99MOt zE5t3^D|6N5U1_}Y?&oBpF>vKLL*fcBbE0qx^w6(D&*OD4OGjmid))B=Ke_w`V(_t_ z^F=wsXj=V;XNM}FyVqjQx!d*(7n1oL>D_D2E0?$y<8SE`v}l{HBsy^n4fse1IY?9y zKPK_N@OnG=q~Hpj`amj0h-5BCCBb=2tWNG5_S08yX6j6XKhzE+g_)K%2kbSh+Dji7 z7C!!9^S!wMt%@LO$_Bn*J?Zmfezc5qn(Ya&$XL<`51YKIs)s96vAJAtnLih|q8Xdb zp)=dJV>64h1v&A7TLY-GoSwPwn1#OR=)g{H;v4torHnQ8pf;#X@45Nn6#<{?4=(w? z^D!*Y_W%3~+W)7GKA~)V>yAAHz9T}AAz)BSzqoE99&M$Em)ZbN5S~dYN|cHj`2G)j Cv&tg? literal 0 HcmV?d00001 diff --git a/src/main.cpp b/src/main.cpp index 3adb73e..6f7675d 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,8 +13,8 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 20; // feel free to change the size of array -const int NPOT = SIZE - 30; // Non-Power-Of-Two +const int SIZE = 1 << 21; // 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]; int *c = new int[SIZE]; @@ -30,7 +30,7 @@ int main(int argc, char* argv[]) { genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; -// printArray(SIZE, a, true); + printArray(SIZE, a, true); // initialize b using StreamCompaction::CPU::scan you implement // We use b for further comparison. Make sure your StreamCompaction::CPU::scan is correct. @@ -39,61 +39,61 @@ int main(int argc, char* argv[]) { printDesc("cpu scan, power-of-two"); StreamCompaction::CPU::scan(SIZE, b, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - // printArray(SIZE, b, true); + printArray(SIZE, b, true); zeroArray(SIZE, c); printDesc("cpu scan, non-power-of-two"); StreamCompaction::CPU::scan(NPOT, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); -// printArray(NPOT, b, true); + printArray(NPOT, b, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); -// printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(SIZE, b, c); /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan */ - onesArray(SIZE, d); +/* onesArray(SIZE, d); printDesc("1s array for finding bugs"); StreamCompaction::Efficient::scan(SIZE, c, d); -// printArray(SIZE, c, true); - + printArray(SIZE, c, true); + */ zeroArray(SIZE, c); printDesc("naive scan, non-power-of-two"); StreamCompaction::Naive::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); -// printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, power-of-two"); StreamCompaction::Efficient::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); -// printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, non-power-of-two"); StreamCompaction::Efficient::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); -// printArray(NPOT, c, true); + printArray(NPOT, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("thrust scan, power-of-two"); StreamCompaction::Thrust::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); -// printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("thrust scan, non-power-of-two"); StreamCompaction::Thrust::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); -// printArray(NPOT, c, true); + printArray(NPOT, c, true); printCmpResult(NPOT, b, c); printf("\n"); @@ -105,7 +105,7 @@ int main(int argc, char* argv[]) { genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; -// printArray(SIZE, a, true); + printArray(SIZE, a, true); int count, expectedCount, expectedNPOT; @@ -116,7 +116,7 @@ int main(int argc, char* argv[]) { count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); expectedCount = count; -// printArray(count, b, true); + printArray(count, b, true); printCmpLenResult(count, expectedCount, b, b); zeroArray(SIZE, c); @@ -124,28 +124,28 @@ int main(int argc, char* argv[]) { count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); expectedNPOT = count; - // printArray(count, c, true); + printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); zeroArray(SIZE, c); printDesc("cpu compact with scan"); count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - // printArray(count, c, true); + printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); 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); + printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); system("pause"); // stop Win32 console from closing on exit diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index a5bd5db..b4df5f7 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -59,10 +59,10 @@ namespace StreamCompaction { * @returns the number of elements remaining after compaction. */ int compactWithScan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); - - // then run the scan on it int* scanned = (int*) malloc(sizeof(int) * n); + + timer().startCpuTimer(); + int sum = 0; for (int i = 0; i < n; ++i) { @@ -82,8 +82,9 @@ namespace StreamCompaction { } } - free(scanned); timer().endCpuTimer(); + + free(scanned); return sum; } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 80c2ea7..ff22e13 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -4,7 +4,7 @@ #include "common.h" #include "efficient.h" -#define blockSize 512 +#define blockSize 128 int* dev_efficientScanBuf; int* dev_efficientIdata; @@ -62,7 +62,6 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); int nNextHighestPowTwo = 1 << ilog2ceil(n); cudaMalloc((void**)&dev_efficientScanBuf, nNextHighestPowTwo * sizeof(int)); @@ -71,6 +70,8 @@ namespace StreamCompaction { cudaMalloc((void**)&dev_efficientIdata, nNextHighestPowTwo * sizeof(int)); checkCUDAError("cudaMalloc idata failed"); + timer().startGpuTimer(); + cudaMemcpy((void*)dev_efficientIdata, (const void*)idata, nNextHighestPowTwo * sizeof(int), cudaMemcpyHostToDevice); checkCUDAError("cudaMemcpy idata failed"); @@ -97,7 +98,7 @@ namespace StreamCompaction { // now call the downsweep kernel log2n times for (int d = (ilog2ceil(nNextHighestPowTwo) - 1); d >= 0; --d) { - // copy all the data to make sure everythings in place + // copy all the data to make sure everything is in place cudaMemcpy((void*)dev_efficientScanBuf, (const void*)dev_efficientIdata, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToDevice); checkCUDAError("cudaMemcpy idata failed"); @@ -113,10 +114,12 @@ namespace StreamCompaction { // shift it and memcpy to out cudaMemcpy(odata, dev_efficientIdata, nNextHighestPowTwo * sizeof(int), cudaMemcpyDeviceToHost); + timer().endGpuTimer(); + cudaFree(dev_efficientScanBuf); cudaFree(dev_efficientIdata); - timer().endGpuTimer(); + } /** @@ -129,7 +132,6 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - timer().startGpuTimer(); int nNextHighestPowTwo = 1 << ilog2ceil(n); cudaMalloc((void**)&dev_efficientBools, nNextHighestPowTwo * sizeof(int)); @@ -148,6 +150,8 @@ namespace StreamCompaction { cudaMemcpy((void*)dev_efficientIdata, (const void*)idata, nNextHighestPowTwo * sizeof(int), cudaMemcpyHostToDevice); checkCUDAError("cudaMemcpy idata failed"); + timer().startGpuTimer(); + // map all of the values to booleans (and pad with zeroes for those values higher than original array limit) StreamCompaction::Common::kernMapToBoolean<< <((nNextHighestPowTwo + blockSize - 1) / blockSize), blockSize >> > (n, nNextHighestPowTwo, dev_efficientBools, dev_efficientIdata); @@ -204,12 +208,13 @@ namespace StreamCompaction { // memcpy to out cudaMemcpy(odata, dev_efficientScanBuf, sizeOfCompactedStream * sizeof(int), cudaMemcpyDeviceToHost); + timer().endGpuTimer(); + // free all our stuff cudaFree(dev_efficientScanBuf); cudaFree(dev_efficientBools); cudaFree(dev_efficientIdata); cudaFree(dev_efficientIndices); - timer().endGpuTimer(); // return the total size of the compacted stream return sizeOfCompactedStream; diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 576738e..b7eeb25 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -3,7 +3,7 @@ #include "common.h" #include "naive.h" -#define blockSize 512 +#define blockSize 64 int* dev_gpuScanBuf; int* dev_idata; @@ -64,7 +64,7 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); int nNextHighestPowTwo = 1 << ilog2ceil(n); @@ -75,6 +75,8 @@ namespace StreamCompaction { cudaMalloc((void**)&dev_idata, nNextHighestPowTwo * sizeof(int)); checkCUDAError("cudaMalloc idata failed"); + timer().startGpuTimer(); + cudaMemcpy((void*)dev_idata, (const void*)idata, nNextHighestPowTwo * sizeof(int), cudaMemcpyHostToDevice); checkCUDAError("cudaMemcpy idata failed"); @@ -93,11 +95,14 @@ namespace StreamCompaction { // shift it and memcpy to out kernShiftScan << <((n + blockSize - 1) / blockSize), blockSize >> > (nNextHighestPowTwo, dev_gpuScanBuf, dev_idata); + cudaMemcpy(odata, dev_gpuScanBuf, nNextHighestPowTwo * sizeof(float), cudaMemcpyDeviceToHost); + timer().endGpuTimer(); + cudaFree(dev_gpuScanBuf); cudaFree(dev_idata); - timer().endGpuTimer(); + } } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 455f5e8..a342b17 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -19,15 +19,16 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - thrust::device_vector dev_idata(idata, idata + n); thrust::device_vector dev_odata(n); - + + timer().startGpuTimer(); + thrust::exclusive_scan(dev_idata.begin(), dev_idata.end(), dev_odata.begin()); + + timer().endGpuTimer(); thrust::copy(dev_odata.begin(), dev_odata.end(), odata); - timer().endGpuTimer(); } } } From 01b9d79677340c27d1c5241b10ea9a34436dff29 Mon Sep 17 00:00:00 2001 From: liamdugan Date: Thu, 20 Sep 2018 23:45:56 -0400 Subject: [PATCH 4/4] quick readme edits --- README.md | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/README.md b/README.md index c80fa78..a4ffd49 100644 --- a/README.md +++ b/README.md @@ -16,7 +16,6 @@ This homework is an introduction to implementing and optimizing parallel algorit * GPU Naive scan implementation * GPU Work efficient scan and stream compaction implementation * Wrapper for Thrust compaction -In this project I implemented the scan and stream compaction algorithms on both the CPU # Questions ### Roughly optimize the block sizes of each of your implementations for minimal run time on your GPU @@ -37,7 +36,7 @@ I believe the main performance bottleneck in my scan code is undoubtedly the mem Additionally, in my work efficient scan algorithm I use ping-pong buffers,but memcpy the data between buffers on every cycle of both the upsweep and the downsweep of the algorithm. Since I do not have to do this in my naive implementation, I believe that is why my work efficient scan is slower. -![](images/Code.png) +![](images/Code.PNG) For the thrust scan, the reason why I believe the power of two array test runs so much slower than the non power of two code is due to some sort of internal thrust library bookkeeping. I believe once a thrust function is called, there is some sort of one-time-only process to initialize thrust specific state. Thus when we call the power-of-two length array thrust scan it has to take that extra set up time, but once we call the non power of two code, all the setup has already been completed and it can run quickly.