From eb640bd9cabc38d0b4c6a026e07abf8e3c1b6c85 Mon Sep 17 00:00:00 2001 From: rms13 Date: Fri, 15 Sep 2017 00:31:20 -0400 Subject: [PATCH 1/8] part 1 --- src/main.cpp | 2 +- stream_compaction/CMakeLists.txt | 2 +- stream_compaction/cpu.cu | 56 +++++++++++++++++++++++++++----- 3 files changed, 50 insertions(+), 10 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 7305641..1f65783 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1 << 4; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int a[SIZE], b[SIZE], c[SIZE]; diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index cdbef77..c8709e7 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_50 ) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 05ce667..1d5178a 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; } /** @@ -17,9 +17,18 @@ namespace StreamCompaction { * For performance analysis, this is supposed to be a simple for loop. * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ + void scan_implementation(int n, int *odata, const int *idata) { + // your actual implementation + odata[0] = 0; + for (int i = 1; i < n; i++) { + odata[i] = odata[i - 1] + idata[i - 1]; + } + } + void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + scan_implementation(n, odata, idata); timer().endCpuTimer(); } @@ -31,8 +40,15 @@ namespace StreamCompaction { int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int next = 0; + for (int i = 0; i < n; i++) { + if (idata[i] != 0) { + odata[next] = idata[i]; + next++; + } + } timer().endCpuTimer(); - return -1; + return next; } /** @@ -43,8 +59,32 @@ namespace StreamCompaction { int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int *mapped = new int[n]; + int *scanned = new int[n]; + int count = 0; + // Map + for (int i = 0; i < n; i++) { + if (idata[i] != 0) { + mapped[i] = 1; + } + else { + mapped[i] = 0; + } + } + + // Scan + scan_implementation(n, scanned, mapped); + + // Scatter + for (int i = 0; i < n; i++) { + if (mapped[i] == 1) { + odata[scanned[i]] = idata[i]; + count++; + } + } + timer().endCpuTimer(); - return -1; + return count; } } } From 069a426b5e71e90c3bacd0deb5ef09c8227f9689 Mon Sep 17 00:00:00 2001 From: rms13 Date: Fri, 15 Sep 2017 22:50:01 -0400 Subject: [PATCH 2/8] part 2 --- src/main.cpp | 2 +- stream_compaction/common.h | 198 +++++++++++++++++++------------------ stream_compaction/cpu.cu | 4 +- stream_compaction/naive.cu | 89 +++++++++++++---- 4 files changed, 175 insertions(+), 118 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 1f65783..7305641 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 4; // feel free to change the size of array +const int SIZE = 1 << 8; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int a[SIZE], b[SIZE], c[SIZE]; diff --git a/stream_compaction/common.h b/stream_compaction/common.h index 55f1b38..d893abd 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -1,18 +1,20 @@ #pragma once -#include -#include - -#include -#include -#include -#include +#include +#include + +#include +#include +#include +#include #include -#include +#include #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) +#define blockSize 128 + /** * Check for CUDA errors; print and exit if there was a problem. */ @@ -37,96 +39,96 @@ namespace StreamCompaction { __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices); - /** - * This class is used for timing the performance - * Uncopyable and unmovable - * - * Adapted from WindyDarian(https://github.com/WindyDarian) - */ - class PerformanceTimer - { - public: - PerformanceTimer() - { - cudaEventCreate(&event_start); - cudaEventCreate(&event_end); - } - - ~PerformanceTimer() - { - cudaEventDestroy(event_start); - cudaEventDestroy(event_end); - } - - void startCpuTimer() - { - if (cpu_timer_started) { throw std::runtime_error("CPU timer already started"); } - cpu_timer_started = true; - - time_start_cpu = std::chrono::high_resolution_clock::now(); - } - - void endCpuTimer() - { - time_end_cpu = std::chrono::high_resolution_clock::now(); - - if (!cpu_timer_started) { throw std::runtime_error("CPU timer not started"); } - - std::chrono::duration duro = time_end_cpu - time_start_cpu; - prev_elapsed_time_cpu_milliseconds = - static_cast(duro.count()); - - cpu_timer_started = false; - } - - void startGpuTimer() - { - if (gpu_timer_started) { throw std::runtime_error("GPU timer already started"); } - gpu_timer_started = true; - - cudaEventRecord(event_start); - } - - void endGpuTimer() - { - cudaEventRecord(event_end); - cudaEventSynchronize(event_end); - - if (!gpu_timer_started) { throw std::runtime_error("GPU timer not started"); } - - cudaEventElapsedTime(&prev_elapsed_time_gpu_milliseconds, event_start, event_end); - gpu_timer_started = false; - } - - float getCpuElapsedTimeForPreviousOperation() //noexcept //(damn I need VS 2015 - { - return prev_elapsed_time_cpu_milliseconds; - } - - float getGpuElapsedTimeForPreviousOperation() //noexcept - { - return prev_elapsed_time_gpu_milliseconds; - } - - // remove copy and move functions - PerformanceTimer(const PerformanceTimer&) = delete; - PerformanceTimer(PerformanceTimer&&) = delete; - PerformanceTimer& operator=(const PerformanceTimer&) = delete; - PerformanceTimer& operator=(PerformanceTimer&&) = delete; - - private: - cudaEvent_t event_start = nullptr; - cudaEvent_t event_end = nullptr; - - using time_point_t = std::chrono::high_resolution_clock::time_point; - time_point_t time_start_cpu; - time_point_t time_end_cpu; - - bool cpu_timer_started = false; - bool gpu_timer_started = false; - - float prev_elapsed_time_cpu_milliseconds = 0.f; - float prev_elapsed_time_gpu_milliseconds = 0.f; + /** + * This class is used for timing the performance + * Uncopyable and unmovable + * + * Adapted from WindyDarian(https://github.com/WindyDarian) + */ + class PerformanceTimer + { + public: + PerformanceTimer() + { + cudaEventCreate(&event_start); + cudaEventCreate(&event_end); + } + + ~PerformanceTimer() + { + cudaEventDestroy(event_start); + cudaEventDestroy(event_end); + } + + void startCpuTimer() + { + if (cpu_timer_started) { throw std::runtime_error("CPU timer already started"); } + cpu_timer_started = true; + + time_start_cpu = std::chrono::high_resolution_clock::now(); + } + + void endCpuTimer() + { + time_end_cpu = std::chrono::high_resolution_clock::now(); + + if (!cpu_timer_started) { throw std::runtime_error("CPU timer not started"); } + + std::chrono::duration duro = time_end_cpu - time_start_cpu; + prev_elapsed_time_cpu_milliseconds = + static_cast(duro.count()); + + cpu_timer_started = false; + } + + void startGpuTimer() + { + if (gpu_timer_started) { throw std::runtime_error("GPU timer already started"); } + gpu_timer_started = true; + + cudaEventRecord(event_start); + } + + void endGpuTimer() + { + cudaEventRecord(event_end); + cudaEventSynchronize(event_end); + + if (!gpu_timer_started) { throw std::runtime_error("GPU timer not started"); } + + cudaEventElapsedTime(&prev_elapsed_time_gpu_milliseconds, event_start, event_end); + gpu_timer_started = false; + } + + float getCpuElapsedTimeForPreviousOperation() //noexcept //(damn I need VS 2015 + { + return prev_elapsed_time_cpu_milliseconds; + } + + float getGpuElapsedTimeForPreviousOperation() //noexcept + { + return prev_elapsed_time_gpu_milliseconds; + } + + // remove copy and move functions + PerformanceTimer(const PerformanceTimer&) = delete; + PerformanceTimer(PerformanceTimer&&) = delete; + PerformanceTimer& operator=(const PerformanceTimer&) = delete; + PerformanceTimer& operator=(PerformanceTimer&&) = delete; + + private: + cudaEvent_t event_start = nullptr; + cudaEvent_t event_end = nullptr; + + using time_point_t = std::chrono::high_resolution_clock::time_point; + time_point_t time_start_cpu; + time_point_t time_end_cpu; + + bool cpu_timer_started = false; + bool gpu_timer_started = false; + + float prev_elapsed_time_cpu_milliseconds = 0.f; + float prev_elapsed_time_gpu_milliseconds = 0.f; }; } } diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 1d5178a..59ae236 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -27,7 +27,7 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + // TODO scan_implementation(n, odata, idata); timer().endCpuTimer(); } @@ -39,7 +39,7 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + // TODO int next = 0; for (int i = 0; i < n; i++) { if (idata[i] != 0) { diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 9218f8e..b9681d1 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -4,22 +4,77 @@ #include "naive.h" namespace StreamCompaction { - namespace Naive { - 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(); - } + namespace Naive { + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } + // TODO: __global__ + __global__ void kernScan(int n, const int pow, int *odata, const int *idata) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= n) return; + + odata[index] = (index >= pow) ? idata[index - pow] + idata[index] : idata[index]; + } + + __global__ void kernInToEx(int n, int *odata, const int *idata) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= n) return; + + odata[index] = (index == 0) ? 0 : idata[index - 1]; + } + + /** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ + void scan(int n, int *odata, const int *idata) { + // Create device arrays + int *dev_odata, *dev_idata; + int nsize = n * sizeof(int); + + cudaMalloc((void**)&dev_odata, nsize); + checkCUDAError("cudaMalloc for dev_odata failed!"); + + cudaMalloc((void**)&dev_idata, nsize); + checkCUDAError("cudaMalloc for dev_idata failed!"); + + // Copy device arrays to device + cudaMemcpy(dev_odata, odata, nsize, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy for dev_odata failed!"); + + cudaMemcpy(dev_idata, idata, nsize, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy for dev_idata failed!"); + + // Compute block per grid and thread per block + dim3 numBlocks((n + blockSize - 1) / blockSize); + dim3 numThreads(blockSize); + + timer().startGpuTimer(); + // Naive Scan - Creates inclusive scan output + int levels = ilog2ceil(n); + for (int d = 1; d <= levels; d++) { + int pow = 1 << (d - 1); + kernScan <<>> (n, pow, dev_odata, dev_idata); + checkCUDAError("kernScan failed for level " + levels); + std::swap(dev_odata, dev_idata); + } + + // Convert inclusive scan to exclusive + kernInToEx <<>> (n, dev_odata, dev_idata); + checkCUDAError("kernInToEx failed!"); + + timer().endGpuTimer(); + + // Copy device arrays back to host + cudaMemcpy(odata, dev_odata, nsize, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy (device to host) for dev_odata failed!"); + + // Free memory + cudaFree(dev_odata); + cudaFree(dev_idata); + checkCUDAError("cudaFree failed!"); + } + } } From dff74f823425fd45e0c031858960be98aaf2d883 Mon Sep 17 00:00:00 2001 From: rms13 Date: Sat, 16 Sep 2017 19:40:33 -0400 Subject: [PATCH 3/8] part 3.1 work efficient scan --- src/main.cpp | 2 +- stream_compaction/efficient.cu | 101 ++++++++++++++++++++++++++++++--- 2 files changed, 94 insertions(+), 9 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 7305641..77ce8d7 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1 << 16; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int a[SIZE], b[SIZE], c[SIZE]; diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 36c5ef2..fd5d401 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -5,20 +5,105 @@ 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; + } + + __global__ void kernScan(int n, const int pow, int *odata, const int *idata) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= n) return; + + odata[index] = (index >= pow) ? idata[index - pow] + idata[index] : idata[index]; + } + + __global__ void kernInToEx(int n, int *odata, const int *idata) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= n) return; + + odata[index] = (index == 0) ? 0 : idata[index - 1]; + } + + // Kernel to pad the new array with 0s + __global__ void kernPadWithZeros(const int n, const int nPad, int *dev_data) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= nPad || index < n) return; + + dev_data[index] = 0; + } + + // Up-Sweep Kernel + __global__ void kernUpSweep(const int n, const int pow, const int pow1, int *dev_data) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index % pow1 != pow1 - 1) return; + dev_data[index] += dev_data[index - pow]; + } + + // Down-Sweep Kernel + __global__ void kernDownSweep(const int n, const int pow, const int pow1, int *dev_data) { + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index % pow1 != pow1 - 1) return; + int t = dev_data[index - pow]; + dev_data[index - pow] = dev_data[index]; + dev_data[index] += t; } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); + int nSize = n * sizeof(int); + int nl = ilog2ceil(n); + int nPad = 1 << nl; + int nPadSize = nPad * sizeof(int); + + // Compute blocks per grid and threads per block + dim3 numBlocks((nPad + blockSize - 1) / blockSize); + dim3 numThreads(blockSize); + + int *dev_data; + cudaMalloc((void**)&dev_data, nPadSize); + checkCUDAError("cudaMalloc for dev_data failed!"); + + // Copy device arrays to device + cudaMemcpy(dev_data, idata, nSize, cudaMemcpyHostToDevice); // use a kernel to fill 0s for the remaining indices.. + checkCUDAError("cudaMemcpy for dev_data failed!"); + + // Fill the padded part of dev_data with 0s.. + kernPadWithZeros <<>> (n, nPad, dev_data); + + timer().startGpuTimer(); + // Work Efficient Scan - Creates exclusive scan output + + for (int d = 0; d < nl; d++) { + int pow = 1 << (d); + int pow1 = 1 << (d + 1); + kernUpSweep <<>> (nPad, pow, pow1, dev_data); + checkCUDAError("kernUpSweep failed!"); + } + + //dev_data[nPad - 1] = 0; // set last element to 0 before downsweep.. + int zero = 0; + cudaMemcpy(dev_data + nPad - 1, &zero, sizeof(int), cudaMemcpyHostToDevice); + + for (int d = nl - 1; d >= 0; d--) { + int pow = 1 << (d); + int pow1 = 1 << (d + 1); + kernDownSweep <<>> (nPad, pow, pow1, dev_data); + checkCUDAError("kernDownSweep failed!"); + } + + timer().endGpuTimer(); + + // Copy device arrays back to host + cudaMemcpy(odata, dev_data, nSize, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy (device to host) for odata failed!"); + + // Free memory + cudaFree(dev_data); + checkCUDAError("cudaFree failed!"); } /** From b15c0deefe409d78a58744da70de0ad53be67cce Mon Sep 17 00:00:00 2001 From: rms13 Date: Sun, 17 Sep 2017 10:01:24 -0400 Subject: [PATCH 4/8] part 3.2 --- src/main.cpp | 4 +- stream_compaction/common.cu | 13 +++- stream_compaction/efficient.cu | 118 +++++++++++++++++++++++++-------- 3 files changed, 101 insertions(+), 34 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 77ce8d7..4276bfc 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -56,7 +56,7 @@ int main(int argc, char* argv[]) { 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(NPOT, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); @@ -122,7 +122,7 @@ int main(int argc, char* argv[]) { 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); diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 8fc0211..609db8b 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -23,7 +23,10 @@ namespace StreamCompaction { * which map to 0 will be removed, and elements which map to 1 will be kept. */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { - // TODO + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= n) return; + + bools[index] = idata[index] != 0 ? 1 : 0; } /** @@ -32,8 +35,12 @@ namespace StreamCompaction { */ __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { - // TODO - } + int index = blockIdx.x * blockDim.x + threadIdx.x; + if (index >= n) return; + if (bools[index] == 1) + odata[indices[index]] = idata[index]; + // odata[indices[index]] = bools[index] == 1 ? idata[index] : odata[indices[index]]; + } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index fd5d401..78444cc 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -19,13 +19,6 @@ namespace StreamCompaction { odata[index] = (index >= pow) ? idata[index - pow] + idata[index] : idata[index]; } - __global__ void kernInToEx(int n, int *odata, const int *idata) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - if (index >= n) return; - - odata[index] = (index == 0) ? 0 : idata[index - 1]; - } - // Kernel to pad the new array with 0s __global__ void kernPadWithZeros(const int n, const int nPad, int *dev_data) { int index = blockIdx.x * blockDim.x + threadIdx.x; @@ -50,6 +43,28 @@ namespace StreamCompaction { dev_data[index] += t; } + void scan_implementation(const int nl, const dim3 numBlocks, const dim3 numThreads, + const int nPad, int*dev_data) { + + for (int d = 0; d < nl; d++) { + int pow = 1 << (d); + int pow1 = 1 << (d + 1); + kernUpSweep << > > (nPad, pow, pow1, dev_data); + checkCUDAError("kernUpSweep failed!"); + } + + //dev_data[nPad - 1] = 0; // set last element to 0 before downsweep.. + int zero = 0; + cudaMemcpy(dev_data + nPad - 1, &zero, sizeof(int), cudaMemcpyHostToDevice); + + for (int d = nl - 1; d >= 0; d--) { + int pow = 1 << (d); + int pow1 = 1 << (d + 1); + kernDownSweep << > > (nPad, pow, pow1, dev_data); + checkCUDAError("kernDownSweep failed!"); + } + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ @@ -72,28 +87,12 @@ namespace StreamCompaction { checkCUDAError("cudaMemcpy for dev_data failed!"); // Fill the padded part of dev_data with 0s.. - kernPadWithZeros <<>> (n, nPad, dev_data); + kernPadWithZeros <<>> (n, nPad, dev_data); timer().startGpuTimer(); // Work Efficient Scan - Creates exclusive scan output - for (int d = 0; d < nl; d++) { - int pow = 1 << (d); - int pow1 = 1 << (d + 1); - kernUpSweep <<>> (nPad, pow, pow1, dev_data); - checkCUDAError("kernUpSweep failed!"); - } - - //dev_data[nPad - 1] = 0; // set last element to 0 before downsweep.. - int zero = 0; - cudaMemcpy(dev_data + nPad - 1, &zero, sizeof(int), cudaMemcpyHostToDevice); - - for (int d = nl - 1; d >= 0; d--) { - int pow = 1 << (d); - int pow1 = 1 << (d + 1); - kernDownSweep <<>> (nPad, pow, pow1, dev_data); - checkCUDAError("kernDownSweep failed!"); - } + scan_implementation(nl, numBlocks, numThreads, nPad, dev_data); timer().endGpuTimer(); @@ -116,10 +115,71 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); - return -1; + int nSize = n * sizeof(int); + int nl = ilog2ceil(n); + int nPad = 1 << nl; + int nPadSize = nPad * sizeof(int); + + int *dev_idata, *dev_odata, *dev_bools, *dev_indices; + cudaMalloc((void**)&dev_idata, nSize); + checkCUDAError("cudaMalloc for dev_idata failed!"); + + cudaMalloc((void**)&dev_odata, nSize); + checkCUDAError("cudaMalloc for dev_odata failed!"); + + cudaMalloc((void**)&dev_bools, nSize); + checkCUDAError("cudaMalloc for dev_bools failed!"); + + cudaMalloc((void**)&dev_indices, nPadSize); + checkCUDAError("cudaMalloc for dev_indices failed!"); + + // Copy device arrays to device + cudaMemcpy(dev_idata, idata, nSize, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy for dev_data failed!"); + + dim3 numBlocks((n + blockSize - 1) / blockSize); + dim3 numBlocksPadded((nPad + blockSize - 1) / blockSize); + dim3 numThreads(blockSize); + + timer().startGpuTimer(); + + // Create bools array + StreamCompaction::Common::kernMapToBoolean <<>> (n, dev_bools, dev_idata); + checkCUDAError("cudaMemcpy for kernMapToBoolean failed!"); + + // Copy bools array to indices array - device to device + cudaMemcpy(dev_indices, dev_bools, nSize, cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy for dev_indices failed!"); + // Pad the extended array with 0s + kernPadWithZeros <<>> (n, nPad, dev_indices); + checkCUDAError("cudaMemcpy for kernPadWithZeros failed!"); + + // Work Efficient Scan + scan_implementation(nl, numBlocksPadded, numThreads, nPad, dev_indices); + + // Scatter + StreamCompaction::Common::kernScatter <<>> (n, dev_odata, dev_idata, dev_bools, dev_indices); + checkCUDAError("cudaMemcpy for kernScatter failed!"); + + timer().endGpuTimer(); + + int newSize, indEnd, boolEnd; + cudaMemcpy(&indEnd, dev_indices + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(&boolEnd, dev_bools + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + newSize = indEnd + boolEnd; + //printf("%d",newSize); + + // Copy device arrays back to host + cudaMemcpy(odata, dev_odata, nSize, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy (device to host) for odata failed!"); + + // Free memory + cudaFree(dev_idata); + cudaFree(dev_odata); + cudaFree(dev_bools); + cudaFree(dev_indices); + checkCUDAError("cudaFree failed!"); + return newSize; } } } From 24c4fc06dd27ad4a120980ea003c6a114e99a714 Mon Sep 17 00:00:00 2001 From: rms13 Date: Sun, 17 Sep 2017 14:05:28 -0400 Subject: [PATCH 5/8] part 4 --- stream_compaction/thrust.cu | 23 +++++++++++++---------- 1 file changed, 13 insertions(+), 10 deletions(-) diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 36b732d..1d2e3dd 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -8,21 +8,24 @@ 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. */ 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()); - timer().endGpuTimer(); + thrust::device_vector inDevice(idata, idata + n); + thrust::device_vector outDevice(odata, odata + n); + + timer().startGpuTimer(); + thrust::exclusive_scan(inDevice.begin(), inDevice.end(), outDevice.begin()); + timer().endGpuTimer(); + + thrust::copy(outDevice.begin(), outDevice.end(), odata); } } } From 59715f4d694965ca2727e0d6ac5a7151043625d6 Mon Sep 17 00:00:00 2001 From: rms13 Date: Sun, 17 Sep 2017 15:18:53 -0400 Subject: [PATCH 6/8] part 5 --- src/main.cpp | 2 +- stream_compaction/common.cu | 1 - stream_compaction/efficient.cu | 30 +++++++++++++++--------------- 3 files changed, 16 insertions(+), 17 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 4276bfc..53f5600 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 16; // feel free to change the size of array +const int SIZE = 1 << 24; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int a[SIZE], b[SIZE], c[SIZE]; diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 609db8b..8410082 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -40,7 +40,6 @@ namespace StreamCompaction { if (bools[index] == 1) odata[indices[index]] = idata[index]; - // odata[indices[index]] = bools[index] == 1 ? idata[index] : odata[indices[index]]; } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 78444cc..14a0b15 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,13 +12,6 @@ namespace StreamCompaction { return timer; } - __global__ void kernScan(int n, const int pow, int *odata, const int *idata) { - int index = blockIdx.x * blockDim.x + threadIdx.x; - if (index >= n) return; - - odata[index] = (index >= pow) ? idata[index - pow] + idata[index] : idata[index]; - } - // Kernel to pad the new array with 0s __global__ void kernPadWithZeros(const int n, const int nPad, int *dev_data) { int index = blockIdx.x * blockDim.x + threadIdx.x; @@ -30,17 +23,21 @@ namespace StreamCompaction { // Up-Sweep Kernel __global__ void kernUpSweep(const int n, const int pow, const int pow1, int *dev_data) { int index = blockIdx.x * blockDim.x + threadIdx.x; - if (index % pow1 != pow1 - 1) return; - dev_data[index] += dev_data[index - pow]; + if (index * pow1 >= n) return; + + int idx = (index + 1) * pow1 - 1; + dev_data[idx] += dev_data[idx - pow]; } // Down-Sweep Kernel __global__ void kernDownSweep(const int n, const int pow, const int pow1, int *dev_data) { int index = blockIdx.x * blockDim.x + threadIdx.x; - if (index % pow1 != pow1 - 1) return; - int t = dev_data[index - pow]; - dev_data[index - pow] = dev_data[index]; - dev_data[index] += t; + if (index * pow1 >= n) return; + + int idx = (index + 1) * pow1 - 1; + int t = dev_data[idx - pow]; + dev_data[idx - pow] = dev_data[idx]; + dev_data[idx] += t; } void scan_implementation(const int nl, const dim3 numBlocks, const dim3 numThreads, @@ -49,18 +46,21 @@ namespace StreamCompaction { for (int d = 0; d < nl; d++) { int pow = 1 << (d); int pow1 = 1 << (d + 1); - kernUpSweep << > > (nPad, pow, pow1, dev_data); + dim3 nB((nPad / pow1 + blockSize - 1) / blockSize); + kernUpSweep <<>> (nPad, pow, pow1, dev_data); checkCUDAError("kernUpSweep failed!"); } //dev_data[nPad - 1] = 0; // set last element to 0 before downsweep.. int zero = 0; cudaMemcpy(dev_data + nPad - 1, &zero, sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy failed!"); for (int d = nl - 1; d >= 0; d--) { int pow = 1 << (d); int pow1 = 1 << (d + 1); - kernDownSweep << > > (nPad, pow, pow1, dev_data); + dim3 nB((nPad / pow1 + blockSize - 1) / blockSize); + kernDownSweep <<>> (nPad, pow, pow1, dev_data); checkCUDAError("kernDownSweep failed!"); } } From a583501c9337d7d7bdfe1e9d53edbcc1cdf7d38d Mon Sep 17 00:00:00 2001 From: rms13 Date: Tue, 19 Sep 2017 19:42:02 -0400 Subject: [PATCH 7/8] readme --- README.md | 137 +++++++++++++++++++++++++++++++++++-- img/compaction.png | Bin 0 -> 18365 bytes img/scan.png | Bin 0 -> 19405 bytes img/scan_1.png | Bin 0 -> 23217 bytes img/scan_2.png | Bin 0 -> 21195 bytes src/main.cpp | 2 +- src/testing_helpers.hpp | 4 +- stream_compaction/common.h | 2 +- 8 files changed, 136 insertions(+), 9 deletions(-) create mode 100644 img/compaction.png create mode 100644 img/scan.png create mode 100644 img/scan_1.png create mode 100644 img/scan_2.png diff --git a/README.md b/README.md index b71c458..0ce2907 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,138 @@ CUDA Stream Compaction **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2** -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Rishabh Shah +* Tested on: Windows 10, i7-6700HQ @ 2.6GHz 16GB, GTX 960M 4096MB (Laptop) -### (TODO: Your README) +## Overview -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +In this project, I explored parallel algorithms to find the sum of elements of an array. Inclusive and Exclusive scan algorithms for CPU and GPU are implemented. These algorithms can be used in a number of applications like stream compaction, path tracing, etc. As a part of the project, I also implemented stream compaction using exclusive scan on the GPU. +The performance of the GPU implementation is very high compared to the CPU implementation when the input array is very large. + +## Functionalities Implemented + +* Required functionalities + * Part 1: CPU Scan and Stream Compaction + * Part 2: Naive GPU Scan + * Part 3: Work-Efficient GPU Scan and Stream Compaction + * Part 4: Thrust's GPU Scan +* Part 5: Optimizing work-efficient GPU implementation + +## Performance Analysis + +The GPU implementations are compared with a CPU implementation and Thrust's implementation. At first glance on the plots, it seems that GPU implementations are faster than CPU. But when we take a closer look, it becomes clear that GPU implementations are slower upto array sizes of 216 and becomes faster for longer arrays. This must be due to the calculation overhead with the parallel algorithms, compared to simple summation on the CPU. + +A same trend can be seen in stream compaction. Both, CPU and GPU methods, take a considerably longer to perform compaction compared to only scan, suggesting that scan is not the only heavy task performed during compaction, and the rest of the algorithm also contributes significantly to the total time. Yet, similar to scan, GPU becomes more efficient only for array sizes greater than 216. + +###### Optimizing work-efficient GPU scan + +Much of the performance gain for work-efficient GPU algorithm is due to the fact that the number of threads at each level reduces to half, so more the number of levels, faster it is compared to naive GPU scan and CPU scan. Mathematically, for array of size N, there are log2N levels. So, there would be N*log2N threads without this optimization. By not creating extra threads, we will only create threads equal to the number of leaves at each level of the binary tree, which is equal to 2log2N+1. Thus the number of threads required to compute the result is much less for very large array sizes. + +###### Performance Bottlenecks + +The GPU implementations seem to be very slow compared to thrust's implementation, probably because I am not using shared memory which is way faster than global memory. For CPU implementation, memory access should not be an issue, but it just becomes increasingly slower as it is a single threaded sequential process. + + +##### Exclusive Scan + +| Array Size (exp) | CPU | GPU (Naive) | GPU (Efficient) | Thrust | +|---------------------:|-------:|---------------:|-------------------:|--------:| +|8 |0.00079 |0.036256 |0.112896 |0.020512| +|10 |0.001975 |0.043136 |0.12544 |0.021248| +|12 |0.007507 |0.05152 |0.160096 |0.02768| +|14 |0.030025 |0.077568 |0.235488 |0.057216| +|16 |0.117333 |0.14128 |0.271616 |0.228| +|18 |0.860445 |0.418816 |0.434656 |0.220896| +|20 |2.12425 |3.16147 |1.45034 |0.389504| +|22 |7.8724 |13.7757 |5.36339 |1.076| +|24 |34.0334 |60.7282 |20.7226 |3.8527| + +*Plot: exclusive scan for entire data* + +![](img/scan.png) + +*Plot: exclusive scan for 28 to 218 array sizes* + +![](img/scan_1.png) + +*Plot: exclusive scan for 218 to 224 array sizes* + +![](img/scan_2.png) + +##### Stream Compaction + +| Array Size (exp) | CPU (without scan) | CPU (with scan) | GPU (Efficient) | +|-----------------:|-------------------:|----------------:|----------------:| +|8 |0.00079 |0.002766 |0.162496| +|10 |0.003161 |0.005531 |0.1424| +|12 |0.012642 |0.033975 |0.16608| +|14 |0.046617 |0.112593 |0.304896| +|16 |0.181729 |0.436939 |0.318752| +|18 |0.764839 |2.9361 |0.541664| +|20 |3.37936 |11.3189 |1.98714| +|22 |13.412 |49.1414 |7.40618| +|24 |52.4042 |189.27 |28.9966| + +![](img/compaction.png) + + + +_Console output for 16777216 (224) elements in the input array._ +``` +**************** +** SCAN TESTS ** +**************** + [ 7 42 17 29 31 49 9 9 40 35 8 48 15 ... 23 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 33.9812ms (std::chrono Measured) + [ 0 7 49 66 95 126 175 184 193 233 268 276 324 ... 410894172 410894195 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 36.3137ms (std::chrono Measured) + [ 0 7 49 66 95 126 175 184 193 233 268 276 324 ... 410894088 410894124 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 60.7475ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== + elapsed time: 60.0793ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 20.6467ms (CUDA Measured) + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 20.7787ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 3.81834ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 3.69171ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 1 2 2 3 0 0 0 2 3 3 1 2 1 ... 2 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 73.5653ms (std::chrono Measured) + [ 1 2 2 3 2 3 3 1 2 1 3 1 1 ... 2 2 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 54.5126ms (std::chrono Measured) + [ 1 2 2 3 2 3 3 1 2 1 3 1 1 ... 3 2 ] + passed +==== cpu compact with scan ==== + elapsed time: 283.844ms (std::chrono Measured) + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 29.1447ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 29.2029ms (CUDA Measured) + passed +Press any key to continue . . . +``` + +##### Changes in CMakeLists +OPTIONS -arch=sm_50 diff --git a/img/compaction.png b/img/compaction.png new file mode 100644 index 0000000000000000000000000000000000000000..bdc1c4122bf271a16b2ab2190f57f2014bcc5f66 GIT binary patch literal 18365 zcmbV!bzIZy-!~qEV*nmK7HQBm2|-##K|+uaBqzcMC8Q+=C<@Y&6X_Ny8C?TKa^R5e z9x&R$7;MCIZQ$>m`+hz5>k0qh<@Sy1x<2{-)B+x=Dbk%_JV8T4L#K5At|kr5VP_he zKd$}x2l$4sErW-KM)8f(-P_tP2bV@pYHJe{$H~U7Gx5cZ{B1u;@b2W3A zJFK_8azn-<-W1~4qNFD3{!YnJjg`|zU&_}e)r}`Gkb(7e;GJ8C52gKj<#RFQ?(5Xa z!?*99{!8N(FYCQmM~@!W)Yj!K011K~z38+okAF@Y`7 zc+O%jQufw?0DqDphiGVSK`EctIl!MemETur;K#v#NoW4c6(7_Q@cExB|2)2b<)6p@ zx$=Mi;#)N9Fh|F=KB#w%*A8iQbu~Rb{dr%$jNeJSvY{|GIaVg7DhF)qr%#`*DTY4I z>utfe8dduzJ)xmFeHuQ{(edSJLtt@n@j_q0Shd^Q>O?&{Mxg4&#A}D8fU{gWdHOdd zRF<72T^IWcZAU6oQ&T^FWOQ@atQT&=I>;nKdXnvjYQ6VRaPsEyV>C2}ej+@G%OQb* z4mX+S;}6r)-@SYHey5Bq+--N;E!#LjE$c^iPEMqRZh>*a-xP|Rz6IYR^g-=Ls|Os_ zz%gP)Ecnl(g&CScSr+E!BTi_;-{j<^;>(+QC-fS3$V8P3t&i|B$|hl`n=Vy50+!xrzqft>1n7M{VxF_bTFD-{?M$EJ4&m$uv z)6&wSrNS8vP@ZlkCMoMN!$rfND=YQsWO1LhvebIIySa63hRf|Si85ujBfLC3;Zler z%p`W*+{afgiFaO&mvU8OUEcgP#TcTn8;$O4k|R!}|Ap0V=@g&nupI z#F|=Mtj(%#rRDqey5;aO0+t)2L(;V6xWSK^SR5!aF);z!u|0g4#z(0aG3n3d^*$yZiw31M+rn%YMk2qXAfBqC|os5c#Y9?|&=ROw^ zp!q&+9Xki;B`hMMRiiA})YP=*$UsBG34?wHTokwNzZ|d?3FpxjV7dyVaeP!?Qb|0! zXLa=jY`316;aQ)YHSJLG?c28~24VU~dU`xWJ5fd%1HGB%?IgkrUy|mThs)_HDJD@3 z5^z$nbfHO7&waQAsfc9|A?%;zK|^yrmZaH=Yl{_WZ*NZ&!6Zm=HFR}#iMT^~*-m%k zoZJ~SRUD$VXBtD8o!gE&wIJqcM z2f7KZ_1XcP($Unsn&-AWEK!kyG$nwW zuU)%_Esc(ju5ww5FjyL{=!!!X4=n`*1UMvJfEedyXY=s#MxY%GI4m+`qmA-zFotYP zS?>MvUqPSa(`lw-Vq}aOsqBH5De5s_dgZ*(+w!QhYG7cXGfMf0N9a&;6AN~r6ZnQY zHj=~msDi=q%y;kJJ<8E(Xlhd7`Rs2{Hed~f6Qa!`on|{w+M&DLOB^sxj`T%o*%vD! zZ=v2hz+mO{6+j_*r52dP5GHZ=F2rHt27YU~YGt}G{i(R!C_hF~Z#Jv^#l#(J>s)EN zif0`YVr(c##w>eow5hr~>z+eYL-_3FF`hN79()UnXq~-58#2_<#Dd%3nL6o=GE|vI zn4X3P_#FyYZ$T+r{+Fxy@1Bb4-KdW4UswJQE|BWO{!cED`r^~8a6KKJ^T0nn>-h|n zD>RfImZ_5y0xWusyR*K&zN;&dTZHOrWji`MkuFPvjg1O%?J)@nHIB1tu{bA5z*RIF zou{(^EU1XEu);^G`^77YGPb`xc_9i-M@Ls!SmUJh>mgqA&CK;XyT5+8b= zfpbCl__Su{2&7M+xUtBuPaXqdyuW|I-FNj9_a8hE;OAeD2j?<<`cx2#;>X1%BqYSf zN_47%w{#%n9M)#qVRAY~Mlvt#p!ht4GG716MI258<<=hy?)V2j{BME9e`gAnk^jOa z|Bc&!a>sNQkvJUNXYo{ovU$s68((Rq?v&Z|lOi^&uNvQHcOrT}Yf|yf=)B#h@bcWG zt>8TYMv8P@Kcsr8r4!-joI-uE=dD_^{B@Z6*D~tIYHaa`8(m3JEdyQmC`hW#VCGwz zb627zeF7GtqDc8FxU4ipB?I!n5Sx$B4%Vx3do+ohAk6MKAHo3y=40Oz z%gO~GD8-3TZ;Rr+SpQ+S;~})aWM|@OLHIy8;{)m$Xc^@* zm7)54%KOjfK(nH$6hWJ0Q6S^EfR+MsJ;TK4(-?$@N+QNRVK#|6V->h9?I`;`IXBcD`oK~*3Rh)D$r7{H(mSkbW6&H zfl3y?KURxn_SZg9cd##0ZW{F=>KlTi>!c5E)Nwz7qWjMs4POMi;e(06X8T8TkXQuE zrgc%tu&Ps+Sn5yR+u24IrgTV%(qB{>dnOB+eI^&n!0sK8yr4_juovUW`Fpv|M4oGU zq|PsbqVH*6QC_fmX=hRb9PU3{qmR^}a_zjxK7k0+7JUyq?$nyu`h*>(d^OJG+RGn> zT5MH1kOM@XI@5a8dPYsg0R{OpEI0(6SK6O&Q$>f(G7L(UtM-RfrpR`j-i}{~jkV{T zI!b+E9_piVfQpi;AC;@CuYRoG#G!$@TslEHdmXLQ3at`!!&tgu|M^iQ)A||2U~-mW zsjPx8^$$h2vzbOazVVFa-MWd=c7HQ;#D#)EruE$He317LAQd+8$pl4`U z-R9eci;C6H6O)zo0?x|H%4%d}q(Zr5dd9}KCHf4k(jUUZzp4OdoLAtVef=tax-DjhmAAS-V$76($;0g2^<4@4#%Na3i1e%nW! zdh7J{OPA=Va%Ld2H2{n@xDS0Gs4pzq@8Hn+xqts&1>KuP=dMiZwCxZ%Sj8QZ7*XC` zBkr3*H4dK1HdC57G?3um!t6Ibn+_pA<2?{&gWkY10cgLzq~F_`yB`vRvR41|_h-vf zu|OPn_@3@>+m0aKeV?Mjjo2^0*OYW^P=PY0p8qZEq7YA@)1HE1`+KN{iO+EemfTja zu!aX8+t)DQ2GSu0^)o{M){n&&jQPPT5(4Hx!UT?O+7DHa2K1d!VOjR^qT z9e@VD_xB@9$0s@wEZIMQf2RA2QG)8{DE}rOT)aTpQ2D9w!8x3$Mz;y_->4^Zz`@b^ zJ}VBv1q<0JhyLGOPN3t*6@ zEe;q^#>|-pJjlmOTJh&a(r!~6!^<4<+$f-5vFGl2dFnTtK3*+`IF;<*RxYONX>Z$y z)Y#qhdNVUos-d*|kF6Cvc;?pgerX_p@lL-55NU7tZ|MWRibF6?-r_e~FZRBz6c4q? zxV}%GGGHNXI<%HUrUXKM6MUaX9wHYHjO!_lpbUtVYs2re7}9&}v{?X=b=E&Ht}z+w zx!+-3QnkM+W;s-@O+sF~@$oa%~dsF>?bk$pVXOveIc=)b8( zA<(8@N$Q)BDXOYsxS5);(1Rl$f=!i!OTS-j&5ZbNb|8V*m;37)?WwwI0l*3J;H@FR zeTaps_oA%r>_EZ8yf-U21z0UiLfAuP)UFwz#Bm@en(y&>+urc zpSm|srZ%S+(B0kLGcz+9C;-b_h^ydaKF7Z{{%Yw%@FLKU3=R1&EV_^jO`A@fI1#VL zH3eJ*=Xs_kE_)~^7Z)Z;4k`#j%Ozh1*~i6}8c&~IrQYx!XXGWXjCu5rd|#;W>9bHGxlh=6g_|L9DH-aAD>~kThF#zR}PXL z$J4Eu%c2rt+#r@;9w?es_ab)S~omwgP1vq$5ZMIn$(o3ZY@sKsO zb#?oCdRmU+&>+p$zf>Iq@Xfa`_QPxx+K>MN=TjO(oI7_e=)AxQdiunvSN6b>oxL;S z7!N324jJXo;ea!F8n-A^gNn@0}7GQ#y?eJqQ zt+Y0K$%OFma4Z(9lS3H)>I;1NwRQl94&6sagiLvXjEw8Ds*;k`t`Kx&t2fV}2L%!l zLPA2V*wWHc5DkZ$|3TwJ)17$wN29$Q4y{*zDF~IDpRcB@tSLBTT?#|T5QszusYqQG z%5A5&nVFh)(jiSz7Em1>9osRKQH6Ec5C@0AdGr?%|50@n74$hk*;-zti;7`y+cxG@ z!)49ek(rQ;?R;phY^(Tir8A$9kiBXjRQVyj6xBzcwu}myYGjw?5s&SImX?$V7+2qv zlqBJ?Xq}u#33i2#1A~HKyNYb^?$OcZdQ|rUK|@2M6-IU5gVg=sO0%XaE;pw=^aKL! zyuJC0LnWqq8;}0_^{ZX&&IZVwVew6D-i8k!UhZ?>JDmvM59gxd+kgD{(bmRn#Ysa$ z4Xw@~Lau0C3It-o8`s|3-~STiqB1i%$Ciz!1`x}8A4Nr*nXwNMGBPr0ZGgR(Mo}`vjeZ!Xnou&iQ{V$92P%PAy;zqW!AVz|H14S0%03@A1&k0SGg;%Df z`BLO>st&f$g<3COht;Z~EjQJnnHd>7U2;2J;nx&jzkWT9UR+!RxfS7Mup09R3|7f- zel+gu!_7w&$?CHdZ9NL$!f1_@JHaeKE$bXjUMs)RzSkHi;q=1FT@NY`tm^@ zy@q(v6$FL704$2l@!(XdZeuLLi@x z$A#jy=0FDGB3A;=9>ma}8cw|(+rR><1AF?>d;g{BAaGgF$XMr>pD2S%Jike?8y8Pf zy}j5Gxmk+)u3e%A#$A*sm1j@*@7)0eFdX;-8|uxg5{P;)%UQ^eb;YZ>vl66f$~Jqfn%yt~6!sttPqM@LwPgGF1=c z>6i=$+{(_%Omw>xGC)sZvePi@}UV9Pq9h;w|?FMS?cp z1OBOcYjRf9S#G*Pg;5kHJHax#A9$qRux9-Y3zl~-<5V0cH2l6{$in>m!IZU4+cUFc z%DaIyqkEeu4^uJ-ho?g6IuSN;vZe?A(DF|&HBzU(#VW+8Q=?i&I=9bhQgj#B&Vb(rEa(fvcd45~M#2Ck^@CJ;)lfd<+MNt}{h#7-C; z06hT@m9*~}K$+-0_IKrD-*k*<0!jOga?1cPkzaKLIhw#TzoWjDw^P9I*zNl5qXo}# zO3H((vip%_%mD-Zf7V%2L(frR7(0k&|3&^=Q*pK7sKuHW@dtp0QwUh9uj57U9?9moJh@U92jN$x$q6|9_vk;1sqEW?B=l&DOAvM`S zI)V1+;A^RPK@z@D=YNXexBoayvs!L1vG~gEPV;7NCYFbBPerwM9=*N2J=GNE@gw|u zaBwg`1X8A-2a=_2&CTPk7ioN&+DA-jhq>N(u8F@okI<4U$1ijsb4p~^u~Q(7Ug*qF zNpy6QG*VR!cBsh96S#Wy^_U3_4J)6Z0Tf8&p%Q%Ogp+{!&-;eEBZ}CDM_%RjJq09T z0Rfw?%m)(_69%Q$KP^Rtvnnekd3cm6dd-Uj3`%orYm>&zz+2dea&B9r!8E}tdj^g; zf4Ck9<#lzrj)EfN79EpV|I$WlYwOZb8DC#3q7!jG2sTG&ps$|~aHhSkOJVv9aZiba zAIwj`^>{Qjf^Q|a)&U||VCJ>=SehgIG6`>T@7}!|;^GWO-aOGQdw(CId4k1mE-X0; z_B*Le;<1>5o-+NNmr6TXkpFFv?-hB7=B44T3(=_uWRY&kx>+qQd-EYBPft43VS1sq zqt#oJ4M96QJ0QIg7#s|$5N8qqOhQ1QrTTfGqyXH}wlHAI2ap3|z09>3Uha7-V#i++rO^{MF#)D+}Yx`b2AdIm6wwZUjC&qoig^AI^LlM|~4ms7~A zf&pmK0ERRs6wrM(NX5`N4{Ih!iLNesY{V!~f_2QjD?l*O)wEWy%{GmmdVe;E zy7ft!eqk>`Oj683XM0sE&P`VNRIZ~Om6TUOs6X=-qGAZr6yz=k7D$b^d*ThLl0z@u zsp(SHiy(Xa9R}2Aw-{6I?w+`r%W3>{pIEzojKGj|B5E##lH2bnX*nD`LgS34L_E8N zcT2sp9F7p6{xK|&rD={#%0bHKRl^+ZE_v^`Wi>3)2051Xef_JO;%cs&fD}z9#SN?I zg|qR8rgUX8Pe)w)L_NVD>kGw$h{XK34N8hh?s5Fk10sV1Pe99ULF%)k@#h@@3?;;A z5kApYY(GKU&^u6Be*Cf?k3CNSj+U)TTC+Q!dD;QSNL!Ik*~#6x;d2oCY&!SPwJGLE z2s-qdJND*XhgZ5W6gfP5YapwMA=Cx;qEaW*DthWP3(vFJ=nTq(M_weE84hG|$=GT_ z-eouCv--~gx>h48%=kb9vgyRpQvqmzG(L8p>@=Ix3X|8)DqC%(();No5=h8*81QDX z-Q;z{yd{lL`E?z$8%-*dp_l!PR{P=_~_^}2)rFUC%P|+x?Us5 z<^v~2iUi&cJorP`g~#F?^b4-{$O`ZB7y98mij3a!0l9Bbcp*+9SLfJ`0u!3_qly9x z2C_>C>cjbyYUH}5k;CL&_wgzmWiNl;1(cHbJ6voo8&pd;r>?GVNM!It*SMt%bK$cs z9gop(Kf6r83~T*ZJ6w}TeS)7p?Y4xxI;hByZH*IxgPzP4*Dx97Ew($Nkk7v*DozoX zleJql&lI^NL3isYK;nt$4&4oqn&PpTD=ZULYrOZV`$M$jT6HOj^k%6&4Z_?EpYYOK z-pP3thg*Iy?G;El?lpjtETQiDR94lS#RmtyH?Y?J(OzU}&^}_^bVT~_R1aG;mT-h< zE8`VHd7@qf#{N80j>q1O%M*p}6~`sk*5d-DIz+=~XTp)vyMsH(VM5R=zJ+7fnYT7u zQYf5$FD!D&K=ybzGO7ho7cI0^I?Y%`mg{#g7>Jv_uR$&^+iItU#qQ3&(u$#+{-Pg9 z@8DmdL1Ip|)2M;(#!Rfqa(x!MR-DnjiBMSeoc6>lL;B7F8K$>gxVOYg$=>+*V;%aX zx~8IRCP@vBi^q1%3-%|pDuHk|e$~~EoPJ2aC%fx+RDhI4@Cmw>P735K&Sdvp$d92>) z=SF3)XJ)J~rW>Bf;=7`lWqo$fT5;qh^F#!as;-wo(TO}~jIG0g{2f3e>W%dV^utmM zQ?DPCMfskF#fv_q(&+<@YZwbi&VJ4>79ICa-93fm z04>rQ8k%K}^v{+Pe!sT8FG(tV)CYZT&0c4e)Jy}$A}5Msq$A-8*t!}cH{;#yk5Sl@ zY*3j;y*M)Mi$Pl371n0=joW(0dg4z+eofs4+dSQo`zQm{h_rvO3Pvv1TmotJNTOp* z;Ha#{@GDMO36|EtmPyhEwQRpoM|krp7&;0;*LY)hHO(og{qhrl5kc-gZgeXzBD#}O z)G^>w-hT2ESci5!z3j!-BI0)zwy~;q(TA*9dgTm$hV5s`l5eTp{L`}cb7$#;B#@Za zbR82~?tdaaxAWwus?3IS)k1Vn46sY2&8V=Nyk!QRWm3ZqvdzuwR@~I+U3wreK6Rp& zZ!=VZM1w+H!eZRFM_H|s(TB7408a2wAGW_~BwO`-s)v?qJe-Gkl$bwBf&PC!jzM0P z?|uqA*1Yiqz`MYE)#}x8%8voNI-ey#P#A*`B7c!zt#}j1uj87aOqP6DE{Z+4d`m)A z0eC_8FITYDp`c(K2`_k|?&`sET%r=VHXf;f+4D?rQ5PLz7HEi>- zWaOC4C+k9?YxxCuaYwZCW43- zUKju-Z1EyR>rW?>oCi1FyZ{&~r|M0W?qgDe;6x&`nyzi(uaI}V8T=eV7+sgGaWChs z!)^t`80FY$WZNOil24FP=SuwAwd71d?1A1nu2O4YVzuCE9xwW7fggKE$0V-IY^CxN5P+L#5Hc_>#<&0>5eu@kKBMRNy2Y0eFZPs?ydTHSvUvIWv%TeP% zUTzi5L)w+A+TGJh%3c*blySMh#_fx4Ig+P(_s^#L6wmYL!!zL#5oe;5^^~8wVAIK6 zU;X~hWc2{sC0$X$QE=GgXWN41?uYm^xU0*~H+c0Gv9}alafClEvl0Yi^S}I@R@zsy z)|Q8%jb0X`4WSs9?gs{j`BB)mFRk6f=jF35(5jTh9{U_YOD{`&@)ld>jmCYW?inUW>Fsg8>f?ydVBKZvJPT9_`lY*kDtG8{Ty<4C{yK?M#>&j^R&L{J3-2|js z6}fo)?HVm=H;9&x-w}H86WZ)82ZjEZf%^r`6M*s{z|YG6(hW)MicDH!V?#RKG>T?) zd}fQ@?vHz>#VfVEC_U5Q@`8g8Bk1{}lO+~-+VNJ&A9MF9!t)N~%C-y>&0wIq(4A|2 zDt4aD!l{@b1%UV>=t1BvY;=5tEpEU9`^YLeG&7YOg>GQ0BF<0+D7jE48q|^cAc<+Q zyVJ?Ji#gfx-9f!(b5f6bFD3O9Kj{*1wqPO5MGKJ`gK{b}q}YCquJketk1p+6W0P~E zI--WB@5@u%-K%;lA%oVybROU%HPoj+#dA~x57m~GuE&${vMO7A<*6cYPCrXOUXVr2 zUy5p-v9oHQ%18j>rWo>T3fkM zha57$t-njkavB`DtyhKF!L_Vp?pe~l$l$5nOCbLp>UXBpha$zBMv9DdAOrUigLTHH ztMDr;gM{fsMGg7(!yg~6-W~us9?z_j)M+6@b*uTpKAgar@QMb7?D}RII>1EKISEfi z-bh1^6rPHTsdu|Ct!9@pM-`=&GjN&Y)6spU<$Vvm$Yl)l+SzKCYMs2Bk-P&V38EB{ zt7GRlZfR##tP`!@8o10V)e#-t8Ov(%PJbArSlUlQZC~zerTmqx23i2pIq4SefTkKy z2($S(-!!o;2@EJ9uUAEUMfvu!bnhA0WGlBo^RJl(jygJ@X9=t>A4|)(OM3b)*yM!I zMKR47EduDq_NzH+q~vg=C$Vuk0y7d<>X^uu$r;U{b$+|Xg98TeAA=#qH;os9NH3j{ zk)rmk%rgDwi2-rqqIp(`DhD zIwF{CHU>xljBV~6V(B7?hS;nt1e|YLjxCw!Dxm(-;P|V;)^#WJ8cE(_2za{@%ekEa zmVvRU(L4j_oNUAv;fSlEG7syn9R==w2-W7S;9WUK?0tb#lMf)@4NZnz`U6CCbyh;e zh2qXo;C7Alct9zYs4MRG33d)Sf^X~=h!r-J7w4=&ce^F?*>TI*7kd)iV; z=;`U=s`{T!T_d3c#bMfGykQP2NTZwH8zM6RQSVb6{(~aI5rk4iU zRf>mz(p-PFc&r8=-ULe1QLnopu|ps=TYOtxW|~nkTR};oBD#uk3ZZgCIavAThVe~7 z8{W61P3D&24y(v_(m-G4tPP&o202Zn+T2KUod%F3+V1^u@IYVr+G7>lLX8thRL>y5 zH->@2mz0MK`!{Q(&ciB^KV;6q=hPZs*Sc~|P3+C&H)`m}&?m49VL2TPDrhTmo*NAF z2^(Ol*yt(tb87LcBC85Uc{O%}^J3&xu+19rDK2pXm8W8Us?jdhMdm4ILR;q*vK1Z|5gz;ORu*crt zN3u2F=LR2vZ@n+pxAFAc2F;eBZKSfQ%B|(l_=Hk_OSkLV=evpNiIr)=fWE*|EQd*> zw&Ie0g*OdgfW6R^#ee!c{rjf}T1$&K#@TyCgeoSDSoX6pO_zjlo}7-6pjJl3*z`X# z7{@~L6N`{|u6p|!e>X>2X;NcEJcH~5I)f?^x3$cg7W<>ng35?rdNw~MOY2vLcWm>g znokiEg)}{ULr2zDBZqY|?e=clko_c58avMgvS)r5qFg8k5JM+DIEn-(It1pH*=Ja~JPC>9?l|zDX zL!N$O<;G#yAluf)&pC-!C!yV1UC!NF@d)p{J^ZCEU2lQNr5x|Ix##C=w^UsDy~IaY zf|Pl97a-AjqwQQH8U6MnMuCe#bR@ieau@FD)2C5r8ylOy{`#wWy3RDcjIi6(ybX77 z9(L(AT8$?O!k-Jseo~EbB)5;jQ`C&01@p=VDXIxAD#1c4*P3Nt`M93m-5=)f!a+MgWVk#>=b*-(; z9dIe@lV`_Tw4Ju^+p2|1Cfxk6Q9qjJ?`o@&1z1{j#uV0iKe&0I^5!+$aA7#hmRwXY zI>1Z5{U#>+`tn=8Ut0bN>q>iG{Vn070@j)>R<}W7AA^`Y+6n)_6>{Etak&AgQO#on zZn>dUlZZK1rSFT14?gH*LW{-O3?N}3Wf(TGCdw{57Q4yzfs@dEP7Z7FpeA6i_v8A7 zv$7Jk@$cm7R8`N4S8VLzcN<~Ons**1s9~yV&uC<>r4V&FTAgGjc2C@kb=q21(AMSC z47_!V0E5A>pgnrU?mBWP0|JVx9d;@Y9wbPnkV~`!)7i%VToyXW}NJwR@uhI=VGEmaU}% zx)pV%=;YwPbV-oi5%~uJds(wirfBc#GemJ>R={G97HLkv$_MeaT2y?A__$T_m za?(upG&191waPoQ7Fri}=6BcScID@3_lz^=-1hng%e*9Vwf8wBqdVc3bNb z{i32G5fPF7D=Gz|qN3I#m5(T?RR+b-qyJ-r?L#ZkkD!(g+PyPCE#2AqXtO5MDZfZA zbxwBn?%v+dA|?8Hr`{sDbX;9slLfp$XB+6?-ew1bI!;bbN5&a^QY$JX`1uz>(fpOf zA0U|x+R~r*<{qzWVE80nZ{j_6cvj{`Klnhb(>UL2H+mR{4gli=y#-JxXr?O~Kq8T# z?tJ$0Nzefo1M0S*eY_Wpi8#iqS(urnpL*N#Oy=e469ol@TLNC7Y!1%O#KeTF7s$m2*y&A3j{fIv5xjfZ?B80;p=z1wq4y zf}l;UdGGn*CE8o|za|?%7e5y#CkYOvm-+*b?KoF?2)TKPV19FP#bdmhmc9ynV7=M; zjnw6F87>h%EteZ7W}~@_PK=Cnoa;#OSwg2+;1^d`7>X-3K~bAk`sH<+h0kSWy1OW4 zizDXHkL~8ibK&HTJiQ`?8PEbj+dWCU7?(mKv-+&Ox>hrC6r3jcEun~bi{$XSWoQ9( z0ATaX>iy{OMd!Z20fQ>wWEU)Q#1?U=@G<|PeuU;Z&8`Dy@x z??i{~fTovp0ktIA8n-ob1p3@VIhsS8=5t%rrRq^5N4x)7>S}EJ46xdus<9QM`S9U; z&?^1~F0sD1jHhR^gj)&6PRkY4=XbA@|FX2U4h{=*l}tD^gpf=Sl$6x3;Gg>9fK{&D zJ1b>Lut~1DL{EY<6cg)A#+Lt>mwD0gjWkB3mF;l(Up`EC{dN{o`xdaZ_P^T~eDpm{ zO7yYq8M(Rq*S%?g&X+AVfo-!qymt8&$AfJ4x19(tZ&?FO14FTReM;>Lg#(L2Mz?Bc z>5U%mc`&3yxFUX98Wx#9fWvF5BNL@u1>X}_!sR~UtWxo6k)w-=XY%fMce#aSmi3dP(qmfX}y09F#s_2KdSTZEB}xE^8f0jrxrND+=Nfk zPn=VjOFd|Su^B3Tkjer23~a2er_a%Xi+|n(IzI5lr2g^n8WxR^+Z^}RHFkrIR=ZVH zR0wC4l@UY4N8fvUFB?@D<}=7`e?cN0;%O~fA~=Bf!9T3I9kg*W9$>lz=Yqw9HABFqIURTRlhaf-e+n}>3VzRBJVT&EkD`bWR6Go zg?ZKOh|44OvTr`8rcQ#{##St7N^Xl2SFhk#2OTM+Q%$D;+b;XuFH_n|kaojIs>^DS z@t}{Y4TD+1;eJ)~>k1f_Uzd`qDky;5xIySrmy4DFLnTEPnB|of#9fc65OFXZ((NG@ z*Y3SN#RkUrz>G@(bC*2u$Wyas0A(Gk1DFc)D5gX^`^sgGnlvRCR8Isz-dxQ)Dr1oj}0O z&hAYxc)yQhm`8*0S4!O%0i;5@?1TT?x7oyG>uT0P#M7Ut_dvZ>(th&$Y?61rSFNES zH~0IKdz)cijv^|T zx9jD$eJ#v+_%18XtN7Ood2CoKap3O77HnkvwDepm=^F*@lAwJ*SV0Z4^ zad1FdKv|_+R9Q7_+Sk|DK?kHsQr2*h^m+%%slw2EkLc}9-dG9@Jb?+5P<|K_6VrOs z?R$Hic!5EgnxM+vRs_^*3;QLE5DAy)l>gLbzV`a)Nk4@w{;A~2Bfw7icvrfu&EyoI zKC|-P$6L;{#g;8!Dur%$_2?B3Mp*JCLYa^2iFi85TcL@9(`w9)p5$*8V2OYzU}x zCxVftR(bB}L5wWfdM zfkv`Un^t=9!$f*^ukpH9DGVA9AAWiA^$(?bAP#GBcBt@vO){lDGBHOjx+gG8Ex{d_ zNT?e&8cgv4u=R4Z2lVFg7k=6RTzNg6Et*rg_flTIK7vDqJ$!E~X-{vY{l5Nc2Eteq zuNJ5?qGUN|8q$pbgJMAAN2^>{?lYt`I(39aM@5|)Cg04ggh){Ky!sG~sb1T<*_$)R z{wVtJ6*|&3)>M##}bYgZM1OL_5w?J^H z`21OU*`2G(Vw10c+3%?Upv$mC-6}Y6TfiiLJ$1;pnIvGD=OS_FQk;0KPJ7rW9;5?^HL2M zyl$CO`UKx0^l84DXS_W|pFF(wYpSW8bDuqFv^$bo^u--dK80-@Bc;d?_PXY zuG|*BvxBUv^4P*-fvr4l7Dk!(3ScW@eTP;H)GVdPkcV@vH%sKp#IY2PhY`XnKKW!w9EV z&z4*eh1?Jo4c=3MpeH5V@!BxG?u@l18f8kSrF+NOW(-DG*CO$z{nHLa>*12pQhKTG z#6E!BGPAPK*IUz~5Qs40aWLI8FW>B~#_6(sr~^?qm4oVGEGqcrD^TX2aK^#OPi^Ra zS@gvSxnkw_+38oJljtDx5TMtW;{+Kh+& zMzv0XF??fD&+R*9D&G83NF!*|28g)cp8-vmq0P>n^ZE01oa^#1Pu{f*zMt2WJ|YlV zGoc-pq7DWDi9mpI!Z;8>W7x1I;v>wUlrAGMf!}|ZnoFxHj@bm;}AJuLsxY{sY`8V zyp_AFRdflQPY#1ubWg`5xt;f$ZU!B()9N1T>S3Tu5)>Ff@61J5Ws&*BrkooB%^|>1 znCi*%m3NHdX|XD?5LfV%3+Gxc_}8bj@e2y-$Ufx0k>+K#Hr*1bDfSG?JC@9@285mT z0?aV48uk!Hm_ddz+gohV3Ev)W%R7~szsIz2i1{LZ3k{+~FobO4uV15MdxnWAXS5#t z*h17}W~(~?4QyK2{>~v@`aZ6Aa+4M9z2(D0?xc2aeOVn|3xB9XsD8?;#`unSxz@{~k<{3yAu*V1NG z<m4SQ%$%Kd3JXatC4O*FbL6#XjOQOMHUA|1&Ri?0l4qWXbcuX zfn4D5b$));0eiFYQNKQa9B8=nbSn+p92`XVuYk>#ei_2DRQp<350> zr!Sf9T?6y=>x7X#4UJlx^h1ahw(bM@t$PU0^|-|vQdqL7scBD*0=q51 zH&;Y5?7I0Q+kitcEO{3F-PX$s^>-Zdsa^VTIM=l?5sjv%@b6KxJnG;_3k)^OKYhAy zM28zX#O&@#0I3>TGmwe^mQa}=%wXgjR%jmE0zdkpD|=b0(?=zb*I!a|BOz(uu>W}( z^iXTEhrMwBkhB0=`Q~M3PK|G_^--PSOJ<`Q_ZvJsEoc$Lu>7>NYhq$9QNBlnd`w$w zo7gB|z}{)%+ffXQo~IxLQjnh?dMKXTOK$gHuKsuQ2%3OE0p;I6FGIcZ{}q-}z5f5> ei;tx4vAxy~{dO$7W}1Q-lH_5T81G>!HE literal 0 HcmV?d00001 diff --git a/img/scan.png b/img/scan.png new file mode 100644 index 0000000000000000000000000000000000000000..a85e2ed13535a32a0f9d40f71c72a987dba0e146 GIT binary patch literal 19405 zcmdtKXIN8P*EXu#y>+V~idzt*s5B`O1?f#iYA6Z_QY2CWNR>{2fQo`h69|Nk3ZWMX zy#;d9HbR|h@ zW4iWWS9D~qeWBf7I$Ddj5N29eqou7)aS|O&h^tW2G-r&C6Yum;4GmV)`~8mxOa^~l z(mPvTaXArsFYN7G!K^OrKi)&A|0XknADcYj2kpvp-0$G8$dmuKE4N`_ zLH~B;zaFPu`M1aa>&pLp;Lp7jmyQ&OhmzJ(WiA&{U>)pqW9J%3zthDXTiQ=@iw`beM4cGY};E)tHSXx(C7Y{V{p0e_2 zmAjyX1Zj1=UZ>5Vz-0WypGS^7c}I5HBNO*T1W!<~b&Dg_5|=N(PfJrwcKP}FC=3QG z9%$=0d*sOHcf_fsp>pQyMyS}>3}3|7qM^37c9=5P1aI`*k0Oh%XHc~82o?@CG&CgP zOlrM}Ujk0_Rw^keISiJlu{hbZWN#ym|8`G!zvX!VK|7AXvhb z;oio^iL!3$>Tg4#Tq*KMBaXqLA)lc#hk4}E(r~5gNR|7Ru7H3*d#b|N=;$cn>uJM* zpFf`TdBixQc{NT!miOyz5V?|B+L&iobck< z36yj8?VZ?An^L37!Yb>c4dcr3sH@A01{ZDTnABbtqc3-|R9)&{vrB>3{r%x1* z3dI|DFjO44Of-e9zsb(d?ke1ZnvbD@GCh#L$!~e@sk-4sgbiFJgM2NBfpgipHr-G= zRle3^$*3|!D6YX~<@8wN^tZ%b_efl5c({HY&%~E6Uv?#dDUsJxkOPD@Io{}}b6uDh zFJ6FI#=9&JOKEF&a5(KgNAkb$n*Dx7K|!Il@6)GGA3uIvMrQXfOg%V-OT$SC*Z$bk zQx};THMl?h$z@92uxymsRA6JWHM%RUs~ex>Ge_8$z6zYu=Owbp&6}Og&DO0^JoD9E zc&&!AvNCxR=I+afX`PEdis}i3Tnc3Zj|TIX_ukBsez3EcHk{iP?gFD|$olX5Dk;6* zYXDk{7l9UUVOO!Qa`p4^9UbV{n4D28{4~`~%o6;T34?=1!3x<`7R}*1WU_i{YeP_yC#&a2!zR-nElyHijn|DY*7lRWL5QuWhevsvRSw|`=0}4~`JbwJRdZuZ$ zhi1HvzCN!+`?qfoozi-3d)aeO6%W+-?yT4PkOsA75$lt`8bgeYjeQKgT5YBR;@*Z` zlIQ}y>^W}H0_LGmSJmYYjL5j!!$Czw#n^ZVm`9K2N6(efrI{Jei6;6=xV?Q|d3pJI zq$UNYp`l@dwQP+MG+19>Kf@t=?P2C7=3dIM0i$jPGc&V4vDbjnL%q^tju0Em1ROwo zaA?-~sSiAyx;i@PVlmwpvsDneI=Z@3<6A3Z{tJEA5vcA=ZNj%63(w`!Rjlg%{=SHu zsHkgy!DEpib{P+uG?DVaS2JDj;3NezHp)$u2 zBO{}n;U*kpO8n-{n_VxXBDy?ueC8<=M)@8y?a74L>(NfHgMxCc-!FEd`&PER9mjAZ zes5j2)KA`-(+aNPVBlbW{ra^|5Dtd}&WirbnOMe$+U0)0Ueons)p5Ya_<5@5KrHxg zJ@|hHfd4i$(BjAc$I$TqHt@j8+qZ9D1_YpmAG5Ns4E*?!r6D9N92-nsY_O)LrtB0j z-OZ_EMUR98_48DW=PM@ehQQy z@Acq8XM&iW<-N4d-~af7|Gb^8tuGRZz=H@lOE_UJo(MK6cd;%sB0|Ay^}R-0YwJmR z`iG8}yFn1>869meK0}q1{JY)V-Nq1R=)D4yN-PGWlwoOQ z@$g|UJ(wto>%u+QsIIGPab2Cinc0U<8=!BYZ{Ey0A7J!&{U3r`J{_#@V~dNc>n@lY z8nT={`(C5Rwv>Q|%&DK#@Vb0LmW{8ftHA`|EF|xGH??*sPR^6x-o8B*9~T{co|QFc__NI|VWX0)+S=q)D}cTxg_IT7pxUai zGP}2P1T*!QI_qor7{RVc;G2xg9g`c)Tie=(YY@dQ$D&)gC2SlVmi)$Ya&j&KWeu1e zdi;SAjd0o}d>0PXJBvhUvxf2BxN)CFBJrE4scBeP*a!)x4qrcU{>YJwE=CFrXkXXi zdGnPBsII_H)+WZZurM!=?$1AGeVIxJrwHhcZn`5!0+W`$t=)`(!V*eiHm5!@ojd0Y z0vrf;t~2dG0(VY+d;k7D)z>!s`0-b)1M%m5iAc4!Na2P7X zK+%D(Uq3B1f}1~mnq89NvJb5_?O?b?olS4biW*G3=~XusKtDqnjy<|HU-U6D!eh3C zhN0Pd-Q1kU->3*EBv0oEs#0+?od7>{4w|h&1KQ#9-=Bsk%NyOAzgdinq+;#oX{;7s zPlG*J+DSQG$VCmQH@pO7c~xQa=NUR_uhFZ{c=QYF)i%MD_Tuo7dCxzfFxx#@6CW|E z4}X3~XMdS99b#ONf}w)?pVt)qG8*8m+nGY_yoc}23HcL+%QK{42!;H#cU)A@lodn{ z#75qwPV)1Es2E0cRNV3>EOi~9pXY*GH!+R`^~1-pZgJ{O3OvZCEgk8!H)O>MW^xM? z#`Xdd^B}YVKb81#bfJmWFR3#+B8eX!G?ATlo`dQ>WT9R^qB5u5q;8_`rSUBFVFq=K z^n*O752C3BlnYA8#jsOcC>Zc_;&rppv z*ElvRHsRG7afeTA1_yBq3r>f+4qc~($IZ-~*P`i=OB3aXYNMwTP+&8eCerff;jG1) zUj4mL&(S}r;=f&&9Vev=XX&)?ib-U1kyBcN3 zsq^@mH66moEiASYVXjMC+1~rX4Y@f&VAc)sKWH~x8xp&sBBu}bD7q{v{|;4eBy;9w zN57tn(DRcH)1>{+X27Y4lVx^yQ$pc~!-VS;o_y-)U}Z-eG(2Yv*W{t8s=#nK=ywKu zhC+wf6`J(-zx?wTz2bRM@V4?i%fAP~UZih?$X!7`6~deSJyI7R$FGEUf6L2AQjv_k z$(wXI2}!(%W=muQ{`-e+tQZZvsQj0Z|9SDPl?d$P^CD|!6JJuL{0voetr$edF%Ua_WYeJ-+2H?D53bk$NI;DEvH?D@IWzDVZ7vGm$p}?g-^wYgzo#xz zF3ue0twHY6d)U->5s%Bu&)3w}uJhe-0U%FPll~GrkYjd6nh%e#;Ebk~E}bTfl`e30 zb~awrS}RG?d24YX)pxx$&#;(RJ4Hy>vZ~X@asj}b8J+Q)cgX9@Ysc*0UYR9eu~;eB zY5O0a#z7(j5{hKu36QdlkB@V3aJXDb0qXx#h^9wMX&j>E;>BWt?;%`QuYxSoyYhu{ zTIA3wHS@QAK+`PZ#uMs1a)g+mE-u*6IE{>n<_JDs-h^*fWimik(*II7&6^EK7A7Wf zyFMY7fePmdki9C9XVWwj#fwIc9(nHiRI^&h*-(1QP`vRwX$r-!Yinmmf4sgS`251c z!X;K#%Tq(t^)owsmYpM1y}3=httVU8U3vugc7>hE^E*0CFeLFI?Y&p0`=)+np+$z~8@E0)Q>4Wn|(6%@Qqi=|XUSNl8foOe}I0 ziI{9T7v0+60brI^y345L905pw;IDd`TS-D^=Y)2fo+-q&wX|%{rbU3UMF!5$W*I_W zvXfkzTAe3&jw1M^rK@UcYQ7#`I<4g+wMcgP`s(EIfEV-t0VbWxtYr&FR=O|rAt_s> zB_$V^K2as2`?CaHUj&JcifRgHmzfUwu?#%h_5}FP5=f?IV(hzj@B0^wm*{4>^!4<- zyuC+pNaG!BMTMlKq}Koz9vK<=Hx~YR)IDXZ%i_00b1&l5=3IAT;)#_fej$yFpnxKl zy@?sH-y)Gf-r{6u7XtpgeM^H73~#{Nx^crr7Ajs|UX2X&!}X7Az1Yo4J}fsPG;acl+yFr4am5oEiI_oje0)5}1NCusBJt@EpWRLCzrk`Z z(QaoVH5S(<`Ki0qKu@6MDcJ?U*K)^^`KjYw{yH!Aw%@?vaDayZ(8oclSL!C`XMjTE zj0#dplegAOXMX%up9KNX1|;0f7cY8&DoNN~ zz0!3?Kv3}5Z@-}~*&JdBkS{!^R85Ggkg&Zg}FI?Sc;5i zK|w({)e-`)gmm5)+2@~wDlTi5HLWfYT3T8Z4t6?J!r53vtwMcSE?rufnrft4ME!l& zt_Q);ZT|e*eTuy&KrzE{B{51|oYz^2k5mG}^Ka%+)`4@-&CmdDB5f0X?_L8ib6^HE1E|Oz zyI>;uD^f<@6}+|g)kK3ifno;y3iRp zB4vT#j#nqo(;XKs1imn^khUQYQ?;#iB~@UKpeQ$nLtejrU0H^tMC ziEv#>O_&xa9934vQ~gHZ6<0e~0BSIhFeNE_`ph{3s2#it3NqL|OZ7o?Cp1+^07td4 zG-o`^z$JY@PaEz`x)26}5gGWay@!f)>Z6OUsn9Uvw7xuHh8sE1)n)E8?fsRy!-@ss zUPYU5fC#|tooHZK-I+sLI?@?@v~`sLa1WqkI)uJbBFL77b(5F8K7Ne?G4}b|K?T7^ zH=r&OUTZUhnj~q+V|4ov*uY z4z;;Rw`WHvQKT=xNvvAt2t-#1u@d79$0WH=Ke3n4-om4B!#?Ej1&7fWON6&RKbu*> za<7cU#&n?pVjQlgZqmNmF?YjwmJqZ+K(D;uv;&%^C&Zo+rPcw`5jf^d>u)H% z^1XI0O|nRFi&#fU17!l>o!VP`>@gKW{v-x4PI3QeYz5t*!E}x>4YDm`<~QaPy(~K- z38jqa2_`?Ea!dfu*Cr(;W%qaB7|h@{5sOd1Cn`fK_>l;K($UuvBK<)?MI0@`g|z`Q znenMz&7FeUNGgo^DQw@r$Wq1S({tWbf8J!*Xn=|fo-yOad;%w1QkVh+Q&*#?_U&l< zXW_tk#zN=u^QKgM`kB2HJac0Nj$T{I@$dAto~)^O39j*{4#Os>r9gVr9OSrUwntp8!T`O zt3WCS-b>I@l7h5}9Snc7U*r{J{JT8Ya(XQ)Dm1XV$@&)NeV5$zE>$WnVxheH z-Y=JBvaH1LG))w6%^*)TstQZQF?e4m6*Nq}K@|A&tJ3Dh{_TWCI%I`I;f77btUmfB zHQ#vdb8~VORUyQ30MPcuk`wN^&^t7V)%N6(HOM;5^D;IOafUs4CZ|LSYZ@ISWN53< zr2=iLaW5tOy5oJXcuKQwab{-p&*1)wjo3`;xOTR0PDNMbk$uVU<;?HhjcJ%#z!^4q z6w5dvXp_n4>Q9Nr*;l&QKI16(ge7`%JnceDtS?vX96oZi_X{ zb>$K6XS-%SL52FC|9FU$e@w=jIko(h9UkrWzZ8?@Jp!ftJ8@W;IK*9-kNt zjgg)>z0r+9v{kIZKuY~M0ydtx_Rpx9kQc965e|pWW@d9DsI$Kc3Q^06 ztYD)kDF=);G|^UdT1Qy2VCPBOW!c!WQAzJFnsbEWr$qa|m$dGAa2c7cz#{W!MCs+c z+rM0#z0M^jQx&ox`#Kl=yRc-+yS#~HXWA4iWJZhanc2p$F7G*2k=7;x(cfrW+~4z| zn<>m!?|5_>xB^93o9x4jX%mBk@Xr+PH5pFId*S!~AI-_mm%t9t)|7B)tUt7%s7R6b z3_b59!uGBPcs9{4{sOs^eI}2$f6r}&@4r^BjPxyDf>3M_>Y4QbOZ z>{)Aq_t%TC%b-l#yZA>?#`)tWE(zQ&T z1u3f}nh|~wZ>WCOz-6Qm@xtfN_6*wdo@RX#Oc4kZpuj$^KRI3(XVZpwe5>gdl%}`O z7(SR-C;dHXGFZ`}o@hx+-)8tCGeOgGuEUg?G2z(e*&0#WxDq-yt&=w478=_}zNP;C zv@$JHP=k7&6Wf~|gZWUfHZ)GwximaqlWI-Yf=Q~H zadwjKW98%K`wy&3XufOcy+lzdO_*f254hJ{93TQ(U$!#U84)*6SxR1m#)tjmH>iOI zs{NxqJ0*2sMRW!Dmi+dZip`3P5iTJ-2YP7hKTQi%wxgWAE3V>Yi7?W)s0zAperuU- zbTT-m(=Rlwb(`y2U0toJYS(>uftmTli4zdxG9eKWY<6}q6-Lca(=APZQ(BUz_xAR>ySs-52Y>zg z1&&X)Q%@dwz5xc6B5FTOc{Pv|A z#>Xu*+jtjvBHAm2!aZxP@qbXTA3t6N=aWU99(Q}%+S<8NKGi;VTom-onPz#u4N)2XM~*YgFaC218! zMaBE~@B8`T*5RZ*v`+l)-iag6yASg;T3t$^FdGDnaMNeH^`vf^(eK*Q3``TEr$tKIki?a0s29J#S*Tps<}ax=ms z?mEqR`Bh?`z5BJk5Ls3+OwsM9v|;SF;NYO<-o1M|I#GM~iU-*K`pa#P?5@^vSfrzk zk8dZONiO5Gt|o1=2NTv(;_1mR{C$csgRC2&*M&)Y!-MY@20PFTn{Y?R<&T{B z*483FPobTE;p_bnFeSI;?Y22mGI}249aVgMfAa?;O5Frs$CywMzCZ%O(k%b6t_igO zrSTBHC#Fr*k+VBPw~O)S7AEcYFri@)Zb22)VD$N!LGgSsD0k8Kh- zR-r^a7?vH+Q11&YeDMRZ)7j)3 zOxQ_WKjG0?_t;4g%Hp_PIeuP+`Z}|D;CA;%5S0|m(uS#u*JJYu$Jf0(=U+EEy{?!! zyk|UXBMH4%z9CZW_ETHcQhY$(cCLzk9wqhRHyT>84vv;~CLc))HPA{Q}$n zBP6-sFHgf|W+t5KHS4vEmDRz7H{S}yRwS{+b(@vZwJ?_J88hsaZ45?KB|7eJQjx-w zAK6I~hGeYO5_(UhoE8jVajB*SVstIP24N|g+0Qd0XyZZc2B-XL1}ko1P?$~lU9976 zCWFK-1R|xr&NG^@L9HEqZq;jJ;X@kMA;Tt$=V?Gl)x=wgiJ=7>%f28RKn^w-o>`Zr z`5xYH%-X6ZOm&;D498}-hhJLQ_rFM5j2SS$_3QhZ*Ze|RJ{7A zS=|la-n>1XU)Zk}FDHsm4I1z|OSuo^`|baLLrcZ>E8C+@Y<0cxhj+5LI8lQxSx1>I zsW-dFJD6~8qiX|CAzGm6EBlZNE3&Pa&;Xi1<6X3_HWwbxPFY(2`8r|W1cy&=-#Cku z>l)N2Y$(xmMN=v-6I7MYq})fom-jIIRJb8~FWBjk%%B&J_PV_6Aflw7*3|Kgyfk3S#fj#p|rEoe;1_SGyqpD8$V1%GNH z{$Dpne)&@G4P35r+d_XxXnC*Ed1cR(psdA<9e>gG;5nu~-6TI#ykRJEXEtFF;3RdE z0;&6>3ediHb~36aQVfyHyno*tAlze6X%taeTG=5A5V8DV9K>15XUTzOC8#k6XSp}8 zQdSY_-|q1hO4A$uk_^*#AlKNjG1kxF+Bg~9uWAx*yq#Z=UCYS?mDAfzR4aoLX-7+1 zA3)*`idAz2^5ZYzZ8HtM(lo1uGl`Mv&9j5v@A*+2zL()uZ0iOYMq{u3r%^sbc)8k#J@Oo2zj7MvUiKxxfi50Z}y}Qo?%2+$gV;Q z1Za?gFD>Pr>cY%;i|9+ijEY z;GPfFm%7whI$Vvc^e1B4yE9exzUEZ7M$^tnG9w6Op(t?I6+=E}{`%}5GH&=lUv=h$ za71m(TsIAJEjuH`^G&6C;XUy|{!X=zo#aTS;09f$^%qc@(jZNM0hdh856E{+8ZAHO z5%x^Y^P82CxbSEqWoh>_+*!7Rb}C?UML01!d<(%Vae5Pv=n1gEbCTP6RHPyN!ojU9 z!nBvw0Nj~S3X|qQ&ntQFl2G!)a^H9tN_mo3f>RMgTTD_kNvF5ve*-m)^2VvGX0~9$ zSA3uI>AIK-6E@QxvohHRxkRzPLKPRaL%9E2bQyMcL;)OlXm~e89`o)gZ>$^E*?#71 z=#jUP6G8EFl(}+_Hf#-Dr35wf5k~h7bRQa~&baayKL{oS4Y@uQH;iU>!}0BZ)6`8LVTwz@Z880?Vth1$opWg z6_f4(!D}WE_ndz<9Zxw&le-`^DBH|QwX;QqhkC@}GWqovy>~ZA+VQ20z;!Wrm(N4X z*kg5KhYp$&&bN!wcAT%TfWH95FNAgt@Jj%(A@m0@*Q0Aic!WX9>wKNGd~BrI9{YIi zL1oJJhcTCBH`=>yy4n>x67U}%M5r0YY9aa}Pb<6)V766mD`ppv2WDJ0x_z6JT@4TO zFW9)Ba>>s`o0f@aY50riOC5`Jxuy;B){bJ^n{j&{=xdx6Ca&*yU!1vsM(ujkE{!}3 zCLF6P0Wm?Dwhwt=AHZ3YE?!4@5~?!o!L?Gbp0W+{NNtmRpZAnmM~{>js)2;mXRRu< z$>sKy5b#8fnB4VQ?IvlFa$tSmUbw)^r|1rnfNoI6GSikQ+Q!O2@umJz{2t%6SKeo} zL9Q@cmghFvFsENV63aD>QQ_8;`|k5gCmMg}R~JB%i)FNAy*u)ay>ykq#{RWZp@s6o zn+uoN<1-7gAMce+<@Hwb^g1xSD`#D(BuS5*jqndYpGnKDc1qxwK9I~;wxz8*dRb^x zaLsTq0*?Taau>NZ5aviJIM;)b*rwxl^Vs$YugNi|e%E)0?k!l&c;M`X43j_=7EpP@ z`RqOqfb%m(xr0hh?qV6xbiIPwF**#F>(_iAfASs=VmWs8(Ai1>XWPGWs$ezly%oGQ z-P+AqADk?(r|#^82D0o+sT;P_wN*xvj__$gq|Cz*8O1_%9IbOfG|B%C3=5nr_8o@^ zC+^oI#DMd;elnuOL<&ClY|~-(OfkUPyh`A4;x;`;JCN!8VfSJ2|3XiEb82j4)2yxnnR!Q zm9e@@!f<}r`U9F{d2^_Fz>fUgTqc~R=2(VimHq) zfJ=9Qr7^qf1p-A=7>x6ubjkN6$s6tvTULr)X2u4Do9fRH99ii#vppRb+eRnc2dv%a zy{QI8_ZQf~y?b}@`~tahV|Jk@B_>A~Ny#Xnx*kuV8aiPIDP8@@;+_nY+2+ec5DR%E zjOj>{+{5-*ecpl5-y?jsD0##C5@mTpQsy#xYzld;{Sn{1yh()zegy4avwNF1)Yoi^ z0;jTXnfF)v-V8q|7TnFa#Xj0=nlM(!(Enayq@`N$PUZCN9l8THVyzS0%gm1qd>c?4 zP)a1FC>1Qdm9ZY(-+%H9?gamJsqZDS&qGcOQJA0?n5rbd(_Dd8kS7`1H1+9Xf~h+= z%gIhJ%_F$f!IQr8L_;A(HLLMStcGlA=O>VW0XPT>s2mifK%{1lHS_ww1EI$F$Ac~% z4EpnqgQNTMVnoV>JL<|biyi*1T_X!=*f`svGo{}ow5HUxQ)xv_#ni`Trm(7iQv2%u z+Up-hoGDvB;yAo(wi7=#nWM*qs51yijsfKZeLrbUxuwd`7Br|B#+Sq(6RCcr>%KGq z-QwJ)#ga8PR!_7d__G&3pjOl!@nKF(vnxBJo+HydZ>xm=kGRX0Ocdz5w%Jq|?+=S+aWz{jG9|%zjHYAxXw(4#ZTgZ!G8J$SUIs zE2hy=aS)EA?VZ<+#&sjSVul{&Nu2Ir^PeBJ^f@8B_6~T3Ih1u>2mV3CLm!+cY5Ro_ z93C{Fqcg(G&f=~t`1qRpwXp54yPG~r=`P}U;J;vw{2+Mvl3M3=vx)Ia-$)v*v9@+Z z9=l}^yLpizN7ZaKO2mV~6F?`*e&tC{tYkGGbQ-MmZ zRxO@K*v(+@f>mYGel65^eLRidVMc+0*z0>EgK+aFV*)w42prIIHl%#<*!H-K73MbO z;-_XaU_#Wg4NyrC^gY-LqpF8ns*MZ8HhpVXJZc?OEZ|WN5fbm|` zkT3t}8FwvjT4ilKrxR6lcg%(Gb8A<_z@c|8+1Fc9FK+n#M(k4AF~|0d-XM;gZ5OB+ zdgEwzfL?~%WWup2AMG)z)XvMEl|x;fig!k@7nMFSD_Lquw1e8Oz@S{}_qHk&y5xBn z4n{predvLgX+N^n&2uOfPKqkpIe|H+AiLyTsI7rt-CTxg%aJq4;^jVmHq>hWla@!J z#}${qfsR(jae`&_j50Z==RAAy@Jh-P=0mDU}}*GyWdRY76=>QXk|IN*JT0( z<&}t^u`lggmWEFtTuYbkCyF`^om&44@!9sZM4JJpndWC!UF4RYZmm_ZG!f%oZbF>A z6GP5H<=8PqMTtDd36>isSM5Lf64^E=<^Z>7j2c10L!{ySG zn)~i#6;bW&f)Pd+^2YZe(Op+?>@hGAu9n>j{1|Rn=c0GkFMFffd|u6^G6}3;`;V*Z zfvFChlSRx9l*^nom(P!Vsx3_Bc#;^K2Dj#ACmEWNChpA6nL!dw`bz*<>jmGjhF74>9%R%egt#LW2B?pD0v) z&9s^u?ghM*sjTE$eK*I^e*NZYPHrBm)^iwP+Oy8Qj>!k-06FyoCFEC@$s%l(CYJ#E zJR|7ARKQTnShh7=1|KT-`>TW$0Ac8eL%8h530kBQ_Ug5LG+K8LR+`oPcfai$WaRl1 zC;o6MK)71oO0W8*&}y4-GMXx(qB|TBfd3e4c*Mu?e_<=Hg2|7zc;$@!@(Ur z`DV7$g`mV-%Oe-T>|cLr7%T)j5wHD*QP(dvnjm2`s_*FwK@3${z z(dnW`5*B0l>p{)aRO`Eij7rCe;0bH?#GUqbhOv_e=H#5I!s~`|Q^iZ?N99-7)6*bf zS{%gOJcZr7=9hXwi4a+5+qM-8Uac_SjR|7$?pJi)<7$_cvSby{71?J3j{S`l?~G5{ zy*)3xlC-VoZ=P|}%zQ@WiqF`Ir~Z1H!^LF*6_B=~T6F$muv`8{G;~HyY`SO4G^fTT z=hlO-H`n{Lo~k_k648MaC(b4oviTAFeGZgZi_kCi7%Fxwh6!Lt2@BW=x5+Orxj_83 zF)rEV`v?ho4Qj)nf(a+YDk&?Ai;01r;pYJ*#IoB${j8|u*JzPTKK2cNnEj!CS-{<^ zu!BopRRl*_cFa~*YkKYUyQ)ZmKhYw)Db8R_Qp!)GtKwY$!FKMVg1jlYYb-aws>?q& zX$uzfVLt32CwT{k`mhyHmrvQQ16cZe{*1!>Zz;Ba?sL+AsCzoNb*Vn(z~K$q!Kdw@ zXh{V%j^p%Qbbja4imc;!TK&?zOK6CX-|6*_@J0QrUG}Ug9T`~82katt#oY8U<6dhL zjBoIl!?Ng7VUL6`Sq;C^xe9B&o_*Z>NJk?7ZV3q(hMv-y=5pNt_C;P24^~B(7}6;Q zy*RMhoLYs8|7WS(3xBkN&!XB8DY0Yq0!N^<01$y6YQD0@iaEm^#LVw5DSAWYWW#93 zNy$AYo~DiBBZQi@9JIl-*zP&4CiVkS~eD*BbtAc73tHcxNu8cVRcO1K-Qn{ zXX=8%exBKmHY<8x!F#37;qh8(@0YWi{TV+GpzPA(dzm=ypNjO6{v+{;L4I~a+Sz&& zQ`+|9Ucl?ACPZVAb z8IW^rw^%jr-+w1!s;fJ(ut3n2zbrjy4=Hq%2W1z(Zo1a{Y)+Ji{VM^xjO-Ox;2)L@ zh)?J2L4S6N%=q{BiNXDR)9ypgt#S#$X9CjZ^%&Z2g-D zQbA~BZ8No8?!y_^tEow~`L5f~hN3y|?M6%QK3zCJf5X*vW_5I$Z~9=W_4S83)Objp zY~_>g=+ynGrxdff>1K{+kC?W)m+R+O>n=rId9pBuWZU;A$==(oZyxgB+WSOlji@8f zw41FI%7=U7T;J^PH9MzC2Ps1QeJy_tD^Bn!9So(%-7LGZj8Qwm_am@gw<;(>#O-OR zuwq2!9LIEh#8-dEx>Jyb@P`IlBTiP?YdL*O?Td2i6m=nWrp(dG4}XdJu#GDt<7wT& zgFEHbofS5nAGfx)%qo&bjs4s@m^F_Fnn`I4uD5)q_a9zos2g3MIKh@Ox}Md&r)>SK ze0*-CZdvDJ-GIVJXr+J;rHamv{{?9K6`=f9&0vj>23RBJN~Ul*p+Gl0=QwgM>+b^) zz`#-?V@pd*Dc`PMy_&+Qb6_W0cbVnlaTkVf)2aI%VG&#Mp+*}M&bi)~c=O0>L4y0^ zcLE!soNH^#Y_VQ@lOc@?bqAkjt+(o9?C;%l z_gWb{pd=pbi^9F!mnWj9_v=JzBuShHk|gG|nf+g|gI#SAS<(#OL4T_M&Sq*VIa5Jl zqQqlK=;gxlqqvF^5)y_T zh~2bPyz2@lL6@PsySwA|vG-Y7S*x3yHAqUCkK>EHygbkcxU;=oF+N}28)6o5Fg7wH z_NlqoGq?YvE${p6rS5FHcKs(shqei3vF4KJOzy~lYD^uN`?9^)WB z!@n`20lR8oU|`+ir(^DlJ|s|X-`_K}|1*cdSib2oA45X}Z~AmdP}J`0D$7ey_ap+mw(v!5UY=m>@h-HY+WwQr)gJlld_KOu zyMQ)0QUC?GIw zOB~s1yuWv_3=V@x-j%LkHNt`m=7lddr_M&B`vE0T|9SgybL+8=j4fa z?Ou+vx3PJrVCplRJ-2TBJ}*y5Qu2qrc(nt}_Ey5@&$rGto)`N>z*mjhT}Dcw06Au4l|GWgzgkaUAD`-suh7qZjPtED&d1r<*+ny> zOG<8$c8~rjcN+W1DebwOWbAPCq|j4!9)1+gZD(O%un<>Wd(q?IU>R(I9B#O>*X{`^ z#=W-xHreIcoSQT!XKGXwBZ8+Euxn&E$GGb5PeCfFB!PpfLjwX@79JiB*qiC;+~s>m zKh2ailUh1HgFdgxy63y5kxrGp|EpqC3WO+-PKJMHK{*#w-3CQ~y2p z9g>+fyy%hq&SW(3=66OVQ~k(upkI+=G8QybidwfP-7x#g4D)^P@F51Zxj0ZvfBcw& z1$7pVlglI3pm8u>+TGUG)%BpgG9$yPD_sp7v@55r0pkF?yE<>8ZR(82d{6fC=g+^4 z$pelnX!ga@ySlg-85{TK8dNwr#2 zCk5fL1kg|p2#{Z%0FB=M6#}tT2TiD;p%rxQ@^peNA_8l62jt6dPPiSbM~@z9YKBQa zSc1Ca!7f>~#dUUd>FDT)faXe|w8>h043>5M_n+cHqb=NI`f~iW5bgsUvCkNUXr|lB zR6^xOBu|o*Yla3YW_@ifvN}l~NuEN((R%7TK{AMG+h{CDkgTNr6pO$6NN+l)HT*+FZUz8xXY6RxSJW; zsRbHIGMs*=pLDTKbRnLB0fUL340NH1F(G=edm|pvRqtg1qE7G!sRFb{PW#q(^~IR zU+C3}qM{;-g}Y(*GVu9|Zgt2zcV6g~BJa1S$jQjbC3A{X4XgZguCDD{?1o;sZ8|dY(l@ zfL_&_e#IZT+1Zit*Ld>DlhN7$EVMCP0>u31#@pSg`2x zHHel@b_7NQ%uxseDNn7m900o115vzt_v_cMi{~|J#FKi-1vJ3)43+(2f^Tn$8%NS5SaWBLFM6)|8W( zIn&ofr*m{vm?8fbrV9;Nqu9@26P;>_v9AhSExk!2o($~L+_|o+Kkvj}^NhU2$H!L( zEQAp&;|Od0>X>8RShDKw=NHFDIIrHkNhNqh1v=I3 zZ5hCjJVltTDbS?8=$R}~eF1zHLB@5u&A=1^2P{?qv9^khK0SHS@Np1*<;s<)%&xro zwF!XkfG5r7UT@~~&*QM-zP{ZBh43TzlyW(^d-jT!d(moYYSFx!gP?&MNIq=trUEJ| z%6xBmB5uBDl1c#nd<-Se_Wmnl4uF9V0H z!9(qeqWAJ39lnhu;KYLYU0p=jR7xRgj+_GY7ciG5oAy%E8>g0Lh~i^b|R6=SiM6i@2g=Qjj0 zGcxcI4L_Y;fS4gPAKYrm`Uc0A(@Pbwu%Lp1KiFUp20?g@3hqJyT7`O1>ox$0iFJU% zm<(Klz>OQ9^79=Y5izpMI~y~cHIN{MIf8Dcme^2aLDU=z3k#qQ8#wh%5nkNT8;9Sx zh`NAY`s6qw?Jo1^gw+8!-2tK{8>*_lxtC6c>*{s`D*fqGr%Or%$w(%+ z47AOI&Nd$eVs8rKA6a+H(%jNAoqTp(rRCl3GrY&_djCtIoID@2(uMWEIOV+c_V-3^ zr=Vn_{ulv?`8?U5qZ2K_uR7K5fyUiR{I(8#>?cUFfQhq@J|ymd)As0UhKHxELBn{h z6HNWS<4~CawjffnjpEwqH!CJ*?SDu`QjMfO)QIDz9y_s$}D5LTx^0!Qp{QggbjLQ}QqF=XwoS%x|i4Tt1Lj${fA2OMj0ooO3k=L>vtNTYV zQt7fNSnJ3D3A3v%?=tyjcQE$yLLq47^ezLfu(b?Jm>tp*w6z&fVFf7jixi#t5Bd47 zBk-!pH^d;Bjmg+97Lll^DD2a2UYW!~&*91_7_Ttvd@jae>^qHU)D zT}u7Mg$ox9DJ$D4fW1LlA^@K1@(^0Y1=u;L^oiZN^@D^-f9B>^Qdf5XEC9ZSdRekU z0044ju#`X^NofaApGC@V*9&c-zO_6;d8VzV|6`AWU3^4c&JI*85)f^*lO)e^aE!tO z=RH?SJ+?S)OHv(D>jl^Q+Cg;7?rOjhse8k<4qrj?gXq5km~Q9KN5{u!>w^PpV1|Bi zXRfnTJE|Db;4BFgU`bWge?20Rnx_arjdDlUqKI6*dUyRnpAwU)I>rZPTWaLwfH4i8}~fBH{i5A_cw)8wsL0pTdt#9^kWP4NVb*u$t3*F^^ChbZ5@!zrDxQ)ZMKsXFdVA zYanl7Z*JI^NeMW~@YaTWveY4F0OG!4`z{V^TSXJSWz z@T>6c?WW+vAHaz&%+CYUZIS*L2?YpbWucnOUo*eIP%F%kwpEL2r4C3MgUe;UJeZbDTGR2F8CS*NWERg z?wDpT&+zGJh3#KP@(*x#yq1+&_q?m77Tgmztg*LT+qZ1m$=gwQtLjy?%k19I1y|Ve zV5uMgSK!+nB_(=gnO#(h4pthA2NX7)QDEC+>}PUMf*@vcAABzWeEVjqKR-QqJ+MNg z7L8k+oqfj}JvQtDX_aJIp%mcu7odj64bPRTa-%_=qQq@ZCwMxVPn*2G{YWTQ7@5ts zVRmp*Z!YwpE?K6+?CkILH)JI)J>xTDoPL7Yf@(_IXf7&sPa?!Wr`BImisHqf(%UXt zBzIL8Lp(FN%f(FR2>mnHAFs%{s_C(>@R6&kify3zjt12(R8&XZsHpzB@y}o2H{4AKE-ETLPsIl^nw|p-!#vZjpsv_hPtN2AAHE!cP(#i>$i4J; zwemQ? z02ULlOFktF{!K-q13%W;!H*eb@FM{GH~0xV0e&3*+6Sq`PbL~Paf>qw(H!QD5>my$7gfls?jYrwg*ObZOfC5#Y03{EiJ9Zq0+TZHR+Ab z%@_~2#Q|YIAF!H0Bt&|5p`fE&SXlUsKh$rB0N>qQblKy)#m%jc_4V@f^mKQ3H!>O- zs|rHILJZ5QA!n(mD)~3e9zXt|Z`4?Wb%kSmwn)3X`nlSBiF>h5>;_Ue3g=*=^9b90;9|Mu{5Jo?AP%+1w4+U>e~7 z+aa2j^1k;Zty}H%jEsyxeOKS?8z@M4ERo4;t<1}rSy|#9O9mDe7Sawz^|Y6-T;Un@ z-Xm><`OKO5hFuYJCQde}N$+WnxM*rd?T-M#%-~;aBKnFQl3Zg=b(NJvy83jWFvKX% zdt)vtDr$Uu9D*aOV^hT)0$`ETo&+YX=w;$xVf2GW}Stp z=|gA{8r`sXl$vIHd)olxjsEuSv5`XJpx{PCWF#UqoJDfL=Gf-O6gwXwj)RTup;My^ zf!^80Wdj=(AI&N=;`sb|a>b|ZZBIXDpjWVd3#fR)PyuY$eS^J0DOk)Zb7sc!L~~de zP3hWHa~zS;N@{~KuKD8Gvu7C?L@TjTQ<3SM?ChDlgzib@Vuw-p1?7Aob#PkQgX%Ac z?4`#1jFY9LrspOmDXB)**4C}8bCV5`M0P6$uXDlo_xAlpDqPd|naj({QVmQ@Ou(|X zQgCl86gT`1MK!Y9F-uMy7|`3xrh5s*uMFpFZfVipGQj%k(&b_>hW?>!DFYjKD|hhk zFN-#MJW^Er=;!9k83yy3?O=DSPW^B>i;K`)U#xxp6Ujr&ciRyu4Huf9?M&}QLLq7& zML7fp4&ppeZu8oE79X*p@7{Ib*ot4hdbQ&Q)NI`>3D)1bBY81|5UD1mTkdYBqXRFS zUg$5t$MDW?d_TINe4<)5h-QaGQYy0VM;nsri*1zF4))8gjir7qK1Sy@+>@A~go zR9IN`0HIb?wMy;geED)uPdm$du*i;L_g^5xL-37_4LW~yx4JN9eYAD-`llgX4UHSY zF`qtVO50AiBt*EX!b*WON}>LtLdkRZ@$u(75iD=!0v@_?yw`v#wdWa@Yl{nuiRlh& z7#Z==w&3xdtS0=7>%m$z<^{P%81Ll~OqfTaxSREC2V!w?aj(Lw!P4z&8p>%xp5nmT z;I9_&iSyan`@7>IkRBw|0j;OP(V3p)m$e)`g7GOvA|0%)v+myJ;HVuO)MrM`_Q6+I zSMhKN&UIGZ&u{p$ zuQzYs{_}gGv-m`6VHQbmAzt2QC)CPZcQ!HU*s)`q4qGz*8{InU@f~)9#SRc|ZjErK z1!b>CuC9fJUnWxu6S@+y{zK3B7luk*2#wK1L|fZq*fFJaY@xrze5pwtO`T8ITr}Ow zjs!X{OyXU>ckkY%(&*&sZ8UHu{Y_1+96X25-2yh7*STu7sM{lg&eKe+nC zAn-rMki#JG|1pOA`@snA)5ohm30jEQ4QM`n+P9X?&d&b*`*&`gywkL_4)*p-)2hIF z+u)zSeE9<0(l1S{oK+TyYi@0YNJb+Mvt-UgzwejfiEfW#!<^ES8<&9 z>DX*7&E*ss3b-k;*04pan^j_+e}5#pv{Phu;?XRA-YmXf_V6R2fTA^KeZL*SNeq0WH1hZB;KvrY;aG!W+{<4ZHZWWG zCRuRRkc2AiskS^k(LC|kr<_h=kef+=WD*0-TQ@UkwEH6yZKt?uLTb36*zZjp9UTS& zy%vF5vzv#K5%qM_3R@?XHp%~asIyT> zHDL{p$|>Ue^PP^i=d<{4DZ;0Z9ZDrDaH@q9(_tEUIrR9iU%Z_V24~54hS&VpQ+0Ih z3|k@i2<)RqD|4a47wH(*TSt`mx@pnK3KWr0si|yt?sx*%H#Rob-0ZP8I6l*!qP(Q7 zg~%;Uii(=_OmD8$N(A-`xoFGx0XyVF;*b>xB&m0_{ilCWxO6%Yptn3eJKYx=>8ijP zu6{g972u~S8q3aC*Q(|qpvUxHpbH)y9ew1;kvVrm!{IEgY?8OBt*tGv-7;oo=7jk8 zCIuP*pYyqj)JYv49$b_)-WC-V1weKf_$1{Rp4o*3vqh?u_L=$l!I6=?eZjkTeF2IC zh(kt3rluyI<~S8q{3i*!NU4@6Y=x1zsZgyp6kssB;WDk^qNh9^aPJ4=BmhGJ(AhX> zYG~lPbt{4n<2d%?v#`x%uw;_DhDIdFLTa;ecQw$0K2l7w(kS4X`gh4@txpT}a`*T5 z(P%UV`_fl2Vf`rZBys0sIbmLt&n0CArbxhF>HYQOkGtB5U0+`ZA!Se6PDEc`{_vDY z5m4i(lK5}@RYMKe1x=6;_KW}F#Jcu?j7TOT4kPgi35(k~t+{QUed zORUtuVQK@Do5lZ}IT`qzXA^Ogn|tt^@hiOp6dK*V-&)|+wb;?o0pg^1t|I^r=jrIi z7Z-WeqSDv;!J-XA+Qt3+%GS-Sa)a9$Qx3I9zWsz79eml2=fA6@s8#VFZ_(hzO9uz~ zgc>MgJoC$eD9eT*AUDE#`7a#$kL=UM&1&9=hMu;ehTqmR;GXI|RMAC0#{OGvZ)NA& z5XWsg;F&wTe-ziD!(Q0{JLQw`$HXXZ$ISIU^KZk=$|y7Z?RNO?s2=*1M~05R!+wAM zFQ>*5X8ywc`)cx>`)_lhXSkAn=vsnLg>_QWF8iuub+eRA&13xK#Nl!W9ltvr?jzwo zjh55z|CVe@5OMpynlmo)$L@l;i`n=rb_lnBZSr!{iKa;GUFZD7-#)_))AYx!{C*YR zTJX(}*r5yh4Ud2M5Nlg>Xj0OCs*ax9t55BRzSRJ|tGQU>m3gZ8!tV!X=M?1u|Evh* zAP5Ow`X$g7GZ7^NY}u`#-;()>d4H?}{?X#}A144d9E`WPm|D!3`$w8?yWFV|g8!dm z1KvmQR0jeQk~PTm``Ezu|JzZ==HJ#Q5e>|5{Fe9JeeK9b_kXg^IOsgh;mn-wefFPmd|tUP;*6ilph@GM5JDfrO{)l-yO3! zSE#XGo{gh}pAupCqx|6b7<5b(p+@!A&wnp50dEV8J^yyxpW_?e5KV=ooS*v>8g}UA zgYiI4q0E2OEqZxL{g++;4!N=})7v|IO&vF|gAIhJlz<3;em|(xGp4cXfRrC)b=Px%-B}$a8Dh zW7Pj(Pt0YiVtYKqIef@LB#sFNcz}Ptx6^~zL|WoAy9qpU3>@B?AmIVPWPg9Zp`jsZ z&$;~DH+~_ZE<9eDr3&~rswHJ0^26xv_vT|eqdtu3;zeyOt?t&NfS4c%_=1o8#6CzA zcxZqv&JXy?=mN`}pPN%sQi`|?(!lN236PA4T%PtiLKX1++d!F-!AEwO7heH*RQX}P z4%S%z9^|4eaiYtZt&zbBFTn{Ae2&YY)W@A`8*KxhKiVR!&vpVaNQjFYd5*mSDTJtq zNTXz^1QiwIYoNy;u%UNNp53f zPM#)ycgLLJz9qb*&t}CwX}&XEP1#Q%Tm##Q zlY)^#XecY-cdVs})6>}VX0plZguCy?IbfOS>9w9bc|r(q0)jr8-lOU7v3CXlkFG}v z;^>r=S1&*iePx1nYkvBv`SaX5ff?+46c8cpmO>KGhK^8GQq9_ADO=vYw2t7ocFnjW zRq_0CEoJ>jo~yW#I{HEG=DR-A-qX|LPMwFdRa8`TbX)-;tg5OiE-o$s`#CDAu=+*W zxYvOiz3+rcIRe0ui`W!spzfy+SC z{}W2M`rh;O+$L?zV=$OkFJHb3385xl0D5KhGKCx1*GBL&@G9pSkZt_I4|E}_b*@vB zlNade6=jS{9F3bl3ln2~q~tO)GZl}4&HN-Ix_~=Yg0cjg@CSaVZq1m*18T_B!eVPJ z;IgQ_?MOwx;Tn9;pmS(+6ql7XceqOVcs&gm#dUorn-zm??5s^A4a+w{BFhn}h{Itp zl3ZMhOVDr6WMHtppGFk`W@P<=8Sdn)W^;Ufz2h8yKZ&yAcilZ3WE9UV{*IyBAnk_D2u)Mis)Ye7fV5OtyIRsOcR&H^o5OH2%ZGqSOGh z=PWDxYUnYnzoUZ)q#(EikY2i-LRQ`X^18cgwU-@+N}Obu@qqEethGo`{DR|i;E12b zUW0wo8ma{lXip4KvA~M}>D;7vFcFj5(ATg3rdaZwCk(HBe;HmyF;fRN6Gv_f3Xb|LV<^Wtb7;t6C0E&C zUaTX&Yt7&w09A2|nzEucCz*#N;{Ks&zhch;SUw4a)ybD1)(NmaENmD26Dpt5e}v-v zq5zt3ld1m|2C3QpZQE2e9v+$NW{yG(uTvq^3j{b+%6r|?#>S@nC}lH!ILy^a;Fb)| zMqn$X1ua@^ZA|C^HP{yNbLJiJPO$4rAnp*Q9z1yPIxtXs_&vq-{6QuFBBX5Befzw9 zGpHPrz8OgQ8;Edhlmv+}SNfjY%e@ti&8TSKsH$NKl9QwWd+k+5B~K9{J#Vr_F2>}2hGV_xV)XG9Ye%94_0^c1QkOT58`pQ%V&Vp zDxYqu4A=!LD?U2W{A9(w`Tf9MLf*2K8Ph(q-9D=H1Fe%OiHPnsGVwh^bx#fGJOUDK zu)e(`o!yX#Y1j{6_F%PU5|x!8cC607KU3u1V|=9+Ybt#QFwk7WIJt3WfR?9J=rD|tz-AlD4*7;hk6bRw z{fLUn6>vV4#6f|O@$J71Wyl(}D zM^Li&8_4|b+lT3s-y5gukq7#O1XR=kW_w=}uiiH>Xbs2kI_chCO|MmU@{OXH?aDvK z3KVU&%NV_=Y>B(TmbEThNKJEMV@JSZ{ zw=-GOn+l0^{v0%`|9Q@kOS!!2b9OwCm&Q)XP9PUF5y z1~w>jgx9R!4E<0etHkyl4WS?fVC?`-Zwj@6(46{EQD$SYGDHub3X2hSKxWd=iY0)*OP?yDP>*6$411p#4zD&{w&zk_d_%je9%68bJY7)FekkmO&lRU7^aQuz zCDAY)-H}$es^v`M!AoeG_C@CT?QaxT<>%}lkY_2NN3G@0G{EAB=Ju`g7}k$5)W)kM zm0B|=$4u<*9DbVrB+r1hP<;b_?(f}TE|13BR$tv!!*~1i;?~{X4rtKrq*9*Ae{xdC zG7{d(^ALw?z<1o(L?$lr;SEP!e1H^suH6E*Sb7Jp_Ku>EDX+j9!}1v0QtU4JJc*Lr z43;qaG?#6tI5Z3%$tIv1^!D&$fuAd`NK0eBmaehTqCYQqNi6L8w;RUz_?xzlqz;Sy zJ!cI&oB|~(3Lr^Wu!rIhC&qMVT|sYG@IEMT3ly7 zp|n@Ma!ubkLq{cF0(|T^6rxH`;T0+^wTx#xcT?)mztBKeack+a_*p+xJH6bns<}ES z(X;xMlK4M!jd;l;seWm{%P@VJF>u%S(xT>SxVdxIox=qcv}oInZb{RkY=<}fixc~!`^pw zjg}w^&l=EPW!|p+Kw1FdqqALXliF*+!no|_c2DOr>tDI)&%65X3bQs(W{n^2Q1a2n zX6QYi$=0PDQ?Kv%kq_6PtHC!*x4LVCYIDD#=WG*6KN0b_C`KOeB$_flMlH7=h6+!! zPT*^OctnKAk#JJ~UAH7yegECV7vAMXY`XgUtebdn@Tfg%&zrgP&E^g5^dti4p?m)z z*mhbUv#;ynXW7!cee5%emi1d3VYK}k+9SPg9=A{GA0~6a&VxvMdZ)W7x(K9@o>ZL9 zMm$1sKb3+{RW_4v`V5_3eB>A%EpgE8w`sQ+FnbE5~UvDwJJP|2YhDvR~j!`=Qg)%s1w`K>}r7AjZMLasB%JVgI?V%#xBjQqEMBZv}b= z&cM+rY#Q3uLCP#7JBV;*-ci5j@7%doBA$9A>`7Xdt92_Gv=H%bYYz?H%}f0xCRp$a z6Uf$W@EgegR;P=e&zcKmqOV0(KzAeDV5sRF)sYg;+gx3w&^41P^` z8_BOA@zv)5vLw!2XXQ-<=)1af;4fJskJaPw_UDDFhZ)Z2(eYY12E+Z7+0m=hA}<&7 z9?C@OxB*y;|EK{<-M)5htn$UtIro2O+o>`;3rDE76dI9;E+csIVhi6;x*=YlxH`T& ze2i)LHEcQD8}hI&xIq@zZO#Ywk9=0LK>9o_du#G>bTz@}^02#?eBDt##OWmlkechTeozotGX^a$fUZ|j8Sag4{NFMCp7+cQNqTl(!R5sBKTNr$m>*x7 z=GO1dKh_Oxjm*8aDc64~g|qNMSY)ormRfW9K~f%MAL;WtrK+cNZ8X$}#V*WO-AKXc z+d`Due!ra>k>h{=ec{6Tz8c@GYNpygpsZ3}H7BcFX*4g%X=c_=L`^UE6ouZiJN~5T zl+gfNvMpztwGQ zQrJ@FgDb zDTR*#=_|R~Hcy{Ey?_5@XMxlKc?}e+UESR|I5~HBb~eiVuTcFoPxD?3sOh+&CuL-m zqq3{mEhpWx>(#mSNVfY8XPT?R*NQLiehjrOF{vnOb-y{oG^~-B=vvlv6J@F24nRI& zBKfdkhwwXE`CTVGMOIh0Kk&@uE)qO9ci{Q6UxbDIG=acds&AiHlyzT4Mdjhcx5dT9 zcb_D-SkO}0@1^5vMq(`>(rLL9v4xNIkx#Cpkb-H<*{f|1=tvW`iv-=Q3#YcMURajo zbdSA4*Yyhu$J$31YhDnIHpsxT4SxS|<3;en_M74Z6Tc80LxU}uAIZ3+y%nUBOFFZc zT+yhn*#0U8=W}PgUBN?|Pn6VOtDq8F=F#EeN2{W@Q8i8W7Z-L$c^^YJoNC1TMU@Eg zvp+17EQQLCyp=L?Rv)P(d^YuxM!7xxG)svYCPCe)DrYasS(Qgmm|Jr7>&PsT65XGp zC#Pz*&{gbf9yW`f+=ErEaegZgvk4=OkpY9TEZ2> zEd4v#7SA~pofYksS`{4-y=F?UHhsBV((P&+7XaH4&LMnysW~>pg?oe~+rE9E1+qcA zX|)mpwWiMsp{}b+Q!5iYvG4Iw?ZCR#pTG61U-4og!(Yz6AG-n28D+?oAyAh2`$< ztGPvaG8Ne%<>{yjFj|y|M`n!&hjy4J8@{`VFx*iQ0oY05pak}IJ)UiWuIS|l0V*b{;6$@OZ2p~lEIvi z4kux~x1{vir@@KF@ym~jkMEJ*=U#LZuAi*f#a>Fy#WeqvhC*N0n;=Oyk7R9jDU7_0 zd1>I(KJAc>(S&U#CC*W`0mXmJyv2pQ=Rv_wL#i?$LxL0CEjgF? zC1(}IbC~OtRa=0nd`!S~xG2zR=y9g~e1!VoT0d-=6(h^r z5bG?cB`heFXs^Gt>V~<2`Q%&&D6w#Z+$#9 zD(n;$Z`f0>h3fe?j=|TL$XAWe;qm~qx!Y!MX47zUtj`6xg$*hcbj~|wRbC-hF!%}> zo4!+fol}vzHOEc2NVb+XmXt+1Pk0RW?*QtM?$nLv z@?VN-Q#;uRVao$JZ@uz*mDzKACKEfK#Xox(jjq)l_3tTAw|Bs8*6?4MhJPNtPU(7( zH<`a@fntq?jZ-v#G~#KL=_9=iLrDAVLjZ8CdY3`SFpU0)8A*B2*xco1F?GPdU<^A~ z$3Be-dd}&&d~iSCt(l0A1_e3mF9-iHo!YGi5u%c@EjB$NA(Fq&%gbdJ-^W?l!}D(3 z`dvy$jrmU2cApD1w2hO;!gp@~_`E7#!|GSTpv(*ig*XYjE#GfNYD2v?Jj>^)0dh?F z)SR5)6^T)l0EuPzl67Okfwi=NoP<;Q+-U{%FnOTBgGks{Wbhi$dMvy*tlrp{S3X-W za>if^qx1DPAvFp^{QK9cGLl&zaqY~wc3)9!Dq+&IcP`tqtkvd+slz^z4^U^7)ad>EHm?OI)`}r!;Gvj@6U9=~e2^*Alio(jB$y z|8?uB!cmAaccLz={`0#-li^=n!&onaNu6_bz-yCS_*)Kf5x=cjH~;5RQ&4R)ffbMxk7L)XmKcEzxh-nt0=@x!`44& zv8a~y)=pGBvm}N!wT)9+Qy`udIQ1H zJ>nS7tNcXIMm>GQlwJa9)c@oIyGldmU7(iF`uceWx}8v>RyvKr@;5N^IJRxlca380Yhka=?2A6@5G*oJ9fIP6;atl(5D^n)mXUS z2I88OP;kJ|lKqdIblT!}%b9MPwCbny8jg=~$6E_;AXA4PZLEm$sE8f*IRR8sDx#Q6|xxQT4{x8cLNla$Qn)>4bX|HER?V7*2G8AWZo20;egN| z+QInMwNCES%plt{nT#2?@NX()2PfpHVv#-U*q>N{*!rTG?o)RlH^^z^^Wp~ASZZcU zKj*B58(lcr;hEB7{ppNPD+Qo3x`yF{F}WAkMn$Tx_OQklQ8e{che1X{ROAwp_iq0h zqmtxJy%R2lOiC~5j=Fp$b!=yxASvEuaRzKCXo>Od;I>8!tS?j{QY3xgA3mQ z_ZtRG6)nR$dq94*b%(p7PGW~S`|_uVgKaFOPEp`s<5Q*l9KBX9#!PKbr{TDwaR z(ILRzxrZJea^rgR==odI?&(Kjt52LehYC7nE1*) z{%)JjcZ9mao%$E9rrREc1*PGV7HaXzO3!13-Guz3O4fJ3(#~e|r|N-id2O}$%A zsqy&Q4TF{s(1{@R9FSM?v?jUQFAxfaa_o_Cn{%Mz5EG0^w43n!l6sfNI83s;HxZ{_ zZtJ1UWfFXDtYq`UIs$~)j?XEp)U!A1v-W?QX~iHqS-4N{CZ+G=vO)1bZK`;M3%Jyu zUdQ@Bd}Y4KbuD%Ynfr+4fFYL#aV{b6%S_=SK&5(H!CV8rY|JOZT}wH3KE*0p2W^KU zq1=`{SKdp@S(r7bEQ@9FDDSVY0iF$wF2t4HQW(E|r|hK~6F8EwgX~N#8!0XvpduBD zUO)SBWszq=7PsgCEt>fL5l|oDxSxpb^mp}B_ZLl>D_=2%27*wRjq>$fUY4I7kOf>k z@RpP)V6BA@w6n3zRoV16vhW2>f2EO<6X2jq;LXWs6GNq~1rrd4Pl1HB&qYsa#N{t_ z=zU@(5}jUu9we*WkqmO@J~Co>MDVriACg{6Cqt*4x^gAE%+Q7GzPYRxTxPn(ejm|a zckgZix*&T#ltHU_%?XIdQF`DOy4$y_BBjtI3t6<3!uAsg#7d|m43@hvQP@y5`JEV5 z!Za*_S_80%z{pnsxlEfVNPhyi<6QB%h-@=7M2vY-wAfS1) z6kz_hHALSlOMUT8XTN%cBPI`m7a&lWNk-Snuu2I0vZ; z{OqN0ZT;Tzd54Dvd;{9moHQgl&o7w?UDkQ|Aue;cb;2k=by{atqnZQ%cLAWq$n@k?&++zbxBhF{Z@47 zT)d4}&scgHe7W=|0+Um5XY8NT?1o9t8s@x7^al&|j8>G%XbI?o(%S z=S(HYnjTk|_S&EQ)joS$a1OyHzA@a%Fp3jwf9m(hM~6rwR(X!`#PO6dkY;7^DBx>6 zz0J+~*P){$`7P=iey!6YLLO2$9_8^42Q{W7sTt+FZ=2o^?{Vt8j3xZ-nsX2RIP9cM zKcA;t@fUH%CxkD0Sbe)KpJ(Hh;ZRK^5}P1?RDP5_S2ac32SI2SHRZ4Q$5RTu`Js%l zm9~my6I0Ps1*z@)2p@}zMzkKFmI`PGHR#H7j*@ya0U-#v2_1J;h~I zER8(*pkG{nXOW8KApq~4y)52)+Y94qbgjNgJ+*;7+p^$Tfmgchsy*jaDpdINu)2Dt zIe|r*Zjfy1$5>O8XxU1;{Cd#R|I?W12%3?xn88&Q$ph7-6!d3p+5Cd zk_K=X7k?fQ=NkQzV-BmWMaYQV|Zzbmx^*T z0rZ6MrXt@0v(Us^+jvxwWI)%SwQ#b_r-iX2AgTt~0N2zq@53^@R|34opcV35MN`Vj z#=-KJHgJrg$X;)l43|DZqP53?NKr@66AHb|D%KxLq88%{2gy>DskTaXkgje-dZu0R zvx?P07ov;ZCd~@3jE>aFclT>5c=2dL;h$nqn z&*11NX$)g3UbulYJBAxrO9I$v*oQ70u-GYAvw=qdWnnHk#Rj6N-?Hnv=~kwg z+QAo4BOO{#bn&SxD|(F#x^|O!RuB#dA+k1QJJ0w5*DxY3rVvf2Iu~w_+ms7Jo18^< znxUUsH>!Msv!9H6?lE*y-)1EVDtRC``>;;1JfWz=DqPD;C0b!*gg(dG*)jW{(i+e( zKTbh(R7+fV;QHxe>AG5H$1cjcS?J364AcNMY!31m9?8$M-%E~@azVe07t5pg&g`q@ zTTTbF)1a*L703^iU_iAMROLQIx#D$4mgt-+ulZk8&n@k5bkQXv2c8sN!@(TM7UB>S`~Mevw=4Zpb~ZnI~E?(!=Bw z9AtZIGWe)<9oTs^ZgqEV2QsprsIk=_M6?D~GB!oR2I49xrlq(>t>dM`l9F1~A5hN9 zBmw=s?IYZ{;zIfqGyAUf#n!CP5~T~#obQtMUuaxi+Q)Bk=90#w!H$$M`;rp}RD=(5@^Z;mF7P9D^az9y?j6Z~KAH!Wu$3v6(b2GWT_o zi>W4na{0Tg!5tNfY(3kap6&)%Aa^zB(w2P|t&v{&!eXOGA}?$ozPDK5w+|=A0s&`v z)lF5EQ%;ECqx-^b55~~rqj($E{vB?unDFJj;Y}N;=Z|A7lHb2qmjSd6s_n?;Tzh`O zs}y;pLzs)*=V4ar#3hLd%#nexQAqmm!L~9rA zdSd5Ul+s)Qu!^vX(A}~$_gAw$U~!@0pTD5gNkqUq2JOs+vE*fx?R9I3%?02~F0@Qh z+ytZ5MOk#|9i6ywRgvrg8rm!sqv~UOHDr6;=mS@g7%duGCa9_XY@Udt2>n zFSdV(NVR58GL~T81$o&r0ZqY9jY|EKn*glM^<0wj;G9y`um5_>t-B7iWas~&WLA~Z zCi(d?f90iIRaR>8BwWrEM~|oItv@lel}8~#)dg+3uxs7B_hzEWiFi^t#G)Zv*j7I! zjb)GyyJRFy_6vJ;<>bXUf&pPCoDM!$5RH;Vt?96`@svG=1&@9(0+Si)+-DV`d4o_; zV!6!$ROIIie6_>$g>JK)4J$dft0__iH{YbI@zIz@*ZK7o6txO#{DRCES1MM)4hyWO3Uk63$LjcY-VS8Jwn9Q3>j*C5F3_8VTP$CFt$Z{U8+oPUz+{#F{qMPld+ z`nTZ3D?dCwA2cvDR&0^)s>WK&s_vH2dKE+&cbaNZKjhRo!_yC7qJtk+?qy4cZ!EK~7$o1kfGY_<#E}NYo(FI(dIqytLV+Dj^ri+ZXJ!#({Jy*(eHo?`U&)&P$ zsp$@^llHvpz|wrqq;X$Y_1A%QQDhn|v&C7nr}2Na8^mV^!l7b1LT*CP3~V7Fd;O(% z$3T>ff2qE&u@u!4HO^yy%edtD!|TqV4c!{!+6dwef8)CNym@2A!rggLT-B%&W*zAk z{OO&h53qxTT-*d(d z2S3Vr?EM=jvsJ+C0-{%(u%b+&uVgMBm6O^7d@qr4wxvi_Dsy+I7?C?~U_SAK_s!f@ z<$ewb_M&sfI~1T0?#wee&fUJYWm}Rcp=M_DkFM~C(0R9|#l+sLtCc6j(6b5~V)mrc ztEzFQpCH%ZAaHmmEO$+>#fxaDDz&yQ9l2BU_t7WAdM>3LVu0cWxi~G|#pZ1i8q`Kc ztnAh8q)+%yhEwlXOjdY!Zni|lW&TaL`cR5R#d;<9=BIV1v__>_Jm@-MKz+XRih7-0 zUQ}|*yWf9Ku-IZKqkCK4s&YemUYh`VkApC?l66*(=6WD=OpcPBxsvSdTmxM)K4Ed( zI;}^KE8O~4fw=Fy@<{If-B<)B zr)<+$;0r+xd%?m}O=V1sgWR@NJ>ldmdwwWHQ0X;rE8?HKt`2?-*#KQ=@kM}toy$hm zfI+`HJO>yE!cU1Meax?ldD3^utM+Gcx-H0bvD5~^uA=i@`Z5o|}tPl7;O8LI!cp z%%x(WdCwj&?oW-zz=_?RnD3p`2XuO%#d6dwj6o!ibd96YMETcJjjZ|VXnpl}*mC9R zV+35ubsI9-wT?@&x40n$*??~WgEzg&C zt~WgJk_jFLjgLn+qjAOg&p> zf-6AfEbE(w1`M?)Y@he86cmfp1*V_3p4r%(S*q)f({$L)jdYnzf9gD<@0am@_8L*J z_QI8@B++asi`cb%3nirv#X;UPyouoMgr=8L5+F$w;=1n@dSMLMh5Bn0XhohE^XEE} zj=dFEm~z*ssNcqgFgthh-K)W_PQ0w#Jov`x2hQcm85bVA<5V1QCR2)Sdl%+1aTix=bMS1WTI{AjJ% zI$05OQZ{IC5bvlzsKd8>OBhE0rvXUP>tQx_x6tVvDrS*UQSu9%UeqNGYN=V8m%KG> z!Ek6RsdSf$kMd5fK3K!yAhSaUZ76CG3$=^I5sltp9 zt!&Pp^d4TXm@lOljb4w>WX?1>Qg;;{ekbLd)~G@SnE9em#w8t%XELGK@wPY-UW^pX0v*}vCExN7VggYm6y`+k>C@2w3J!QfIL zbMBM;n}aW#kAae&|0{_)GhixSJqzekqBzl-h$51*u7bYmj!fT+l>12-r9h=5>t5a2 zC|<5~U1(an6SIQbO7g%(JMIoOMeX#sddeiNe$R}T-=|IMsJ5V2RbwKUq6_^LiUQfny9?_R_?DwtTItw>x_64cwf ztie>=@#j5L(@9C(z<~x;iJVEc9IX`0MO?Wz|G-ot_UUv9O|m$3jTQ5Qy5Va3_)`lO zkq%$|%XM@JGY%7AI*v5xVo>L_m8rMKDGu+?IQT{%NF)4DJk!wt?7HiJpFUmM>E2Iz zDbV6%F$1%+=EAc-0@QLiMRph1OPuu-}jjApj02vWpT(fg(wb7h`S2Mnk3lch`FC8+l zPVS2LJIcJPRuljugu$+O)wZpPcZHP)e>(Ly70LOb>Ay!#D7Q2CqYu(&e*E}R``)_X zx?=q9$o}qJ2`=$clgO@(T_ttd;^nQmgQ2a>@y|{W(#Ys0xx*=Pl(lV>6X|r`e4|)$ zmsE1FH$76E&hh>z)!nmfelz$Hi<|6By9d^XQ#h5eOZd4tYkz-#x4~!pIYt;#B6RV9 z3v}BvcEYR*`v3axllF&0sxTNDnGQY?fbEKkj<)OjCc|sqKQM4>i}AfY6bf}KpVry^ zHIWa}%3rtSP?D4XbxRH(Hy!q}mTK8ua-y;|Jqs{ERJN63D7E?Z=P%do2+W zbURnOtW2Amn?pPJ!~OO2^qxKYLT0@tB(x5?eJH=fscUJWZTnGRxQ1V;N_t;ROe{{+ zVRLoW_-&v+c_$)Yn)<_06prq9mKf-Wm@ac^)-PG|%FKPOMV3lBe7C}=%-@?@9f z;jIU1QvXCYxc;nsQD6U1-s(;r?tk@X5&d7s#>Rq6Cn&#$9!%T>Q_1_ma!{IsZDj*+ z+r}oCBO>}A3ABu@OTxXK#Zy0aaaki0@EiWpx%$On0!$)cDyHqgH!~;a^gsVRE(X(T z+PQi)+oz|u@sI;>RZQ0+xL2XF@?Nb4cz&rfp%D)CD+RaUkes*cp!)^WuD1g%e(i^< z-?nIDzBk)nq^7o(<*w%k0j7`AP;aBa!h(XxYnP_5Pw43A!0=S%_UX0l%kSlPIAh3i z%9NEBqtS|^3xmb2b#)3$$gHgAzJV^kp2^>U-0naqvs?D$=-%+J7QA|cIIaAzo+?%a zA0G)>Au^21zW`kJ<3jFS3C9V+K48?oLDrf|T_@>*UVU!dh3 z|ASvgP%o{&M<#=@l*#aDaMuKw_Nz60u_x-Pk|_JrSvxRwhc&=T0Va$|??G^AB zyFLNI9e{gUBBlLP7YNFst&U;R{^vU3v}s`c8=R+u!!oaaQCC+N|J55C(>~xnilLMq zoGI~HY@IX&(rUHOBPt399o2O$+oHn4u+*=sQj^Xp^w;@K!K6A2=S3G{T5G zj8t%Gq*d%~4l0DRMEiuwPh=ypw+?)8knZm8js4wa3p2AbYd?KgzEgv#t)pJ1zP@B| zNl&_J0%^65wYRG)Bf<}ikb-&3ZRVpOQ)H9{_MWDeBn76Xii$+Gu+HGxfI_QIJ}4Fh z3^08q;4;-jxy^*s4enzi@9hQo`yTIu%gM>z4lFG#1%5Nhe}4zqVD}!dy|?e)jrX1g z!h3dF%zHg!xmY#+&ex@n7s0@ksPx`?r^hI)ceThjv#3bh+omy!%gD$G+?V342AUQJ z&^F!Y5!sW{8#=1s=A!AjgPms4QG*gk8KTX2b+CGTet!OW=7&;^CM^jPoN8X+5-oxo znB$8V2aKk!eZ32(T4DvbPsXb+Ud#nmP|$sJa<3^#n$qaW>bL!MfcyWc8k#!IPK|l!*2_-5e6iuuF z2|-HeRihvz6ltL(uz%d0v+kU;ANE5&-J6@-w>{5Wp5L>eo=P|pS`*J^v#ljV&X}f9 z_W)-yeNYrSZ>L~)aC#vS(!y|GQv&YYo@K8OQC4#>kIh_*2yuowY9m9X{{X{I)9%T z&Zk4xbco*F+B(yDjvfe*q@p4xmdlvEBm)}%iDF;ncy?1$6XeYX2L?=0s2a1-dnT@q z{zj?1z>90ZciT^}=&Z%cH{N&5wwlSRS?;83jfKnzJ3^`ABha@d?n2~sY3>~n9WzrP z>AJcShN$Q`gqq32$MW_LO}kH7WKR@rD{To>PKG4$112-04+W;{lHr50vIwXScJSMH z7}#mdyZz1XF08<%nM*bb*txkmNSnDvLW~#D(HVj<7@7n=nvlo1e}C)u4RY4v-3A7Q zez;!9xh1)Iy{)PFM;Am#1ALDH^UDX&5%Yq(B$B=CaG?%;E1fbtvJ^b#l$KR4zvuRt zeS4?vg62#I`tV31thBgz%bFG0DeGb%=$?^v3A!F~4t4m7a{P~LS8v-t{@tM;*BtQK zb~f++l|%SU8O^^*z0!EJwYqNcH&f3MkIh+Gb8!J`Qz+veNoFc-&Pg5n{SrW7gN8*c-?v z*s_O+1oORnRX1(QcEun8%nlL|yMkGgh>(zw$Vf)```RaygYVMZ_QmTRg?R32s~Tf=H8la7Dm5yW|VYB*Je%Dhk}O>YlIvolga3fwzoG9?(6KN zoq@qHK0f}_z<>&lgmQ3pKG$;4MTL->m`GWVjuX$a<5d-y#>U2Jrhyjt57%xal1S9l z%HM-Lyu21`_2@f&Owuzmm*xd@94oG8uu?pI1`fOeWbYDIX{rsfr%em4c!lxGa6v)e zzBRCB@LWdE&=B9pB($1?M+suA+k^Hxe*B?N_oGukA3S)YlDp?U+iL0O4=nM-?;_pJ zvbJzqwP_qu;GDWT*Yl}xn5SlgpvdJk!`t8AAEMYWU~6h>-gBacAB9uySbZ`*`uQ$hz&HVJ{1nze2yZrz$cj!8uwn9n&(au`OW?R}`wFqa$A2zKpv{?uQ^a1ZKG$ zPE34AR#sLL%OEiXvLi>kl~q** zc$4w_83K)%hA7aV7aS#cUO@yX4V6=jAs%_<7Z; zRb=AQB7s{$SSYJl?Rc)frl!WFcLKMrXa!<%O7_aI+bN;+Td??`EunS?aOA-;z zik6jSGZ_;TW8#U!(X0FT0sZ?8VrpO>6K!WHvE8btC(y(B`ub-4f@7){*1jpH@@b}y}hUKZHJI|aj%oSJ3G5_7VY zPIn&hW^{26jP$YkKJy^Y4~_{yhT&e};nT)LoW8gP?rC2<{&9sWJv;j~pJu+Z2a~xP zgTYLt#DY-(AJVF^pqidlSWuu3ZOq7MqFSBE+p(T&88v>tLYna1MQy~%N>(lNr;d{;kz@)UcV{YbUSF_nD7xyB6YO}eq zaWQyQ_ze5P_#(YQuv;lr)IK?jir$4MUmV8sZZ~s4 zRnllTZ8n^6bBi47wIUMo>dg`jnjvzIbVQk2Pv1U0t$JjlGPd6A~9k ziw2HdR0@9ETL2RKF>#0UaG?ACG7P~SUs@u`VXbC5M3(MqC)HL%HJ-+$mm*lX#6-F3 zqPf})BFf&$F#GTu+Pb>BApW<`LK?-GMEqED2j;aoR@)iB_v8^z_8WC?Dn|-&77h*$ z9F#0($1VCiX9hpA)-p?E83uHO5JXdx z_B06|C`P99>=Le`-7LkM8rsjsCMJDtwh1*l0^B16X}P_6)P8a+B@#zYl2=nJl7cb| zp?A+&X|ZyVc9OhXy+J#PTn;2d00oSXKPp9%UIMD5Dq=)hMy7c{d5!-F2Npzb-TFZ! z&Zvi&jUyO4HvirquWrgABRY8uq?^QWS~qGh5Fy-Oyc}>~u5z-oqx2iZ`K-nhEZ$j9 z;geT_Uwhuk%zUcYRXs7^G?*J;4kxie!y3a1 zzVs${n3Gsl`OlvwBnW8DLt*aqW2q@Ak1{BM-T75jUH2&dT;(lWJOR@ZA0HMxZ+od+ zC^?^gF2dK4FcexLjHohd>?8rOqSiM4fGQpHWt0;m8in@gf9cMhJM0G-h2t1Z(<&K$ zH+g&1ZehK|{eHQAY=@PZd^tiQ#Cn{f9bNa?K9l19g}FIZg`l(WX$4+6EL2k0{ZdZ& zn!4>@uX#cgrl`))H4=FWeqL|663uex-yM@Wg$+@}Og2oG(lu5CI+kVjYv`o92grh$#{LEzxX;>!lYDNepj(56i^;c@G|FXlj0Vd0obd zT1=(x0K8-nBaP+!u^7pL39PAU+2hCfsHpOaiWAMin{+1f3kfxsrkpkB&ZR@JCWo4u1U)KWy8+eTeCN@klSr*+-0;_nyae zWDnZByN%ufxfdEc%y_xY4K7rauVg&Ux zg`hsZeQTG*@S2$y1Sn~Pltry!Ucc@&^Xcab<0B>n;#TnN^0rz7~xBqOSGrYS2&gj#6g zyZ^yfQfwxT&OHkQ?Thx>goQWwcp+>t%ds>-%)1jHoEx^cdH@r>=;3LpBY|nv)%p4P z)1fU_-7N`7I+5V3>e{y_n=;Q#l*%kJgdWvZ`E d;Zn|=brh17clBgP99&&tcfjeFBAZ{Y{1XcAZIA!} literal 0 HcmV?d00001 diff --git a/img/scan_2.png b/img/scan_2.png new file mode 100644 index 0000000000000000000000000000000000000000..5d70329652d682151578458378ff2e8ffc698bb9 GIT binary patch literal 21195 zcmeFZcU+TCw=aqfR0LFh0s?|clP(}2T?IryKza=kklv)%0E&W0uaT})>AgchI)q-O zO9?&nno!OJ_1*8;d!KX9J)isM^$&z(o_S_g{jN1@jo&MIDWa>CSMl)hh@@XgDBI9*Ric`WOw!&!_;~iOUGZxoK8JQqm0e zqp-v!)UnU|Fm+YMe|9Ttw-l^ppS<_ds@pHxYiB8w95hmQz{GkEp4OhWUjeJxNvxgu z2mJ4axCnm2a35iAf*)Tw&*9-cfr179{$^1YeAom3pT5D{VE}*s$2b4>_}MrA_W1wJ z#Q1JdtFcl`c_t=>gw-Z-NePM8Eco`UaFNtKtdDYqn-cysfvne~xHeH$BIO?(46Sm^shk)X zp!d9X^=jDsShb6trDe`2Eh#A}6%{LgSDXNk=ZU+2Q_@~+r;*FDKRIW+rD>x!9Naj6)Ei59Ce zUmqVD5!Yg07;`6+n8$Ser4lHO+({+^FGfPi?0?vsHk>rNy5Zj3obNCS4R(RCA^=alTSk70@ zsx>6>wp(UaR!ME-^=p?ee-Uz8+g%wROaNSwNh&EWHrWq-qe3gP|MtkXGlq970C2l6 zZ%t}u628!#s6O@OA6s$WRO$fX63W(r1jF^o8pw%mUv^8-9qkI+)T2+-Um8vJ38sEj zI;<2=Nn*I~-Fu22dKxP zp{{O1akh$sipoDYjbB81tRX8a1^vV_D!mCv@bU3w65F~mGg+jIv$NUIRsKHo?IPh}Vdu}EA5iC;z;ZJ%7y_~YcGkkMdq9a(msOxL zb8X!+`t2u4Nl7`{WimfAY|%Y%#_+bODJYGEm5C54gMpS)xB5X&Mif_{k&Ag>U*Fu4 zv+pht^!4{o z5D@0+RNVLMKqBoXkFQ~3tr4L!5mwr@9x!QXzsg`YXgdt(N$MpA+EjGPUU7GegWv?P zh-4f(4?J((xs#LgWaJxXypxm&zs1jA2&o)h=6eq&m_vtAb2f^vUfp$|pr9y!{o{Rp z{zGO(%ZgosJJs|hf-Hww$x^|Yt7$mve`d4)A58w4oA~W9{=*skXJ7K)I)ncg6C2jD zvawwwB*a-)3kwSagXB?76%~3Ra3zHdNQd08n z-Mi{JFlxsz_$PV~7`V($B&5m_Jz1me2xgIxkoYdtdLm#qUtChMnRx-Hvm#+|IKR`H zj$do4EIj5V<48E*uGNI#Zkbw^>Pcp5YU(zLnyM<~5*{9fYs&jYBmp;N;Ex|aE?v4r z=Fb5BF*2qhH)S(ko@C)X?j~~GaeKbAi}s_snwrmt58vYjF*J%HG^5|3kwTX>&%>-Bt%52W#;DQ7cN|I0Xju-m732S#WYK(dte>O z)BNzWG;}Yl*CdYLdJ=#ta&|5HZ<>mV_k`?hZSUW|KX8YQ3>hb2H#$1%JP01%*-?~b z;P9{U9oZ4=!iSB>-i!zf1|}6ABvT6cAx1;9f=B`<|QuL8J&?|2Z zZ5PGR;r>22(Pn5IQ8OpD(d3bga;Ccq%=94z3oEPmvu6W$>;)%gn*7f;*SN~wT(}R7 zD9^sKfwJg65os12NN57Qo*K}NUT zO?YLJ7`Tj3y1~GBTV0#D@x%uhm9t$+xT&t z&w?|isMF?mH2zIqIO1Q~xxQQo{yV~ywT*SlRcofuqv7hm@f(nyfy4Z~``^7hwiwFn zcBkh+PI@Xn3nzVcwy-up<(u%PWRG4o-NunK7mjIsOD{cBxRII;xbrVMb|Tg6Tq$Sh z{?7Jur!uUd`-RT(Go!bERti`HOv4KuwU<;&`Ng)+esdt>cuRRWj%HntuO{dZHKhAb z9uLPQOp94{cFz5gpuR8e4&&R`L2CSQ})I3N_KgyWO? zb028&hUA@^1B1^H80|O{@ip!w6m1e}cK;tnM5AWLDmlQ3eYe)cV zP5aYwbZammkS6w>o%F;!ranA8?CtIK@bIvu)di*j|3h=531h(Qm3_mPFJDSZO7c5B z;N*m3H-Jg_b+`u6i@P%G{2rD802{z?AyjHvyFK5;K}6_S6~7IP?2~V2LpCB10RaI! zyStH8=0n+s8x17s3W*3uv{si8$f&%R%VTxf*35lK+`WthasJTYU`c5yVzZfQcXyYF zh^WvPOy#v;2Y7q*yoxcv^HFSM6a(VSFLs zCtOogvp-#d7QBOk|7;E6s_L3&UVxn2yXYJV8o58|#(D4Fx9sffdZL9pcz7gl{+ius z{!zVIti{Fvf?CAY-rSs(i4qU5TOBP16nS7lK-)V3tPG6=v8(kIV&Vzi=U$lizBb;?5n2O|8_v zJ)8qF3Ya`vVy+5}`i<3VGjno?uFZJ@A3JW)zaj%m2rtjxOKkjLbL|U)dLmM4M~5+7 z+N>}mLwW@n7Z>;G)8*YVGJivmZ2_s1>Sa8<%f6{4_149WhKB^Ydt`J>7cZo7(x!tz zp+6lmB{PKK77}ussB~Di&QZ(SF@#T@WM^hpUd5}s9xepSZZw?y4mFBd%7pn*D^bva<06UUM}Wna@Y) zWH&R*qrS6TN@la1jEu~S7rz(?@UGdgo2Cz%Nd#8$z2$pYhfy*p(2$XpWjJVWZ>QyA!*CKK?2&9j&T>@g~VKJ*J|#1{3f&a<0+k z=$>4L8lBlFZOcgA-6iA|QpWG8?^U3>RgRYJ5zIK{qq%^CP=POK zO8`ZG^us~cg&JsP46k_#r-8nHb08HwCZ-+tK1I8n=a2So`ypc_`@tJkV9cLCf9A6q zFV}DSOhm;ux3FN>cMi{a06`sIQ2sytd(h!)4dD5IdWj zT$F)vCV+%Ds=gI~U72Obx)Gix`rqq;Q}J5`1O?@m(*o<|-=={g6}+u3}8z8}2qScF1%!5Ok?Pzh1N~Btf8Kq%LF;xF983 zTH1=~!CSc9!1P{D}&-cea*hJ z-npEtG+sEtYbY$EK#FaIl}C6M2D!rJ4SB0EE>RJwN81*TbWCsF?xARZ4-^MrO0|fo za!7ICCCFhwh^J&!ybPYHhm1Dg-jTD63x7HWkAngQ`pEGHf91PaE1J7s`yGD@8}v)) z4bQj(T-K(}N$<`sLBYtKQI!Oz_f|3{w(|HcPz)Ap=Wz~ikuYf@p4sn+CEhHJT5mD_FSR?W$+9@(6Mx{ZmKaQW~4JC7QrZ28vW zqtviTfs$rFZPng*c)xRqIfE=0MkQUB>n}W}is^7~_;NU+6vf%y@e@`Vbp|}WXlx)U zEdb>k%6V}bi-R@ZaAO%zNoYgbe$MKQBgMh-9k!Gm!CUai#DInSBA*9R9;M=3$pAve zRIcj^&Um6xn`GeFw(AF^KW#bHEifZ;P)pYSn0CqZprcgmjscsuvjr)Ku* z1W0S<%QIYg3w<>N^QrEh+fA0;C=&GvHDwpIRW^f5_s<9vwsJKj*#iLGsx5yogjlEN zU%WCdqH<9`L!u-$hj64nNhf#_MfPa@l3g+(D!7{Uj7hQpV`EWG#sRluY?csjG1Ib6 z&a!7u?OV9US5`ZJoLao+?6PSTE~ot%hM^>rPsh=fcDC?o-i*?Tz{S}uTdVB{p2PJ^ zoZWAx?ax*febMmal*uhPUyap^<29g!pW_Y9e)c)JSViYX?q4`F`%JiGfTlyF__9+s z$0NioXL0s81nLDZVR}D~GI=1`Fu3Z+!Z9YZ<xo;0VE2MYpEWNyx)JkdFfV zt|mC6oHZqBvagB`nR+Y?zQAp91+2(RiUMB9iw>r`T5Ek|6Hi%6jY@>^RfY1O9gAVY z?->Svmd)sRdWuwAW~a@q%<*COF8zl;)}s0x?sTjtF(3m=2z534)M%)0hDP$-F1-l0 zQAcM`Y5%MV7Y%!3`5kCWJs&m+O% zeoX?L1}8&}hMW!*r4|o0T$XUL=o6j2ususzA9may8;^l)o8#B;l=|~``D`I$-LM{2 z++fr1cYH#g{P_EuN5kLWP84-4=fi|?x^R~R`>u2ClaN|NWaFRty=Tmd*0*$iPYK?4 zt0C&o$WCCT8h=s$>zl~bPMU)BnKF+hr6_}WV`!J<{KvC|h0TcU$j9cg^VTmkWuQI` zQPEoH>Ym?|xQQ#-Hg$3DzAL9s`@opTjibaTq^@A>#A73mQp(Z*fmen_Le7zY7W3Tu z2W4@p29Y*~qq_!TG{b+vXFdf5A?y}{aLMR@ z$jD0HPS^kZ`7*%Wh0Q%SrbyN@X7jUThhv#0X)pPb%r8l1I2$94=kZ;_EDhk0iwUJsH< zNb<^NaM$<>+k0vr-6=8pX#`& zH6KOlwGeBt zX!zpAr^=lEB{ELk4v-Y=nJGMrB7j-~AWbln_cg1m<8z04LzJ@mBug$INhEpR(F^SO z9o49|?sCMN5u;xaIp`iwY$q8hX?RSG>r}1K>RwMz5373a3s**P96(RM0RX%@1kT@_Gy z`JnAJoh^{q&BA3NPmuU{zdj>@)2y25bJNh4p7atr88c*@LRA3&ZDs!rVS4!`WDsJX zh7sgHa&XhMI?dp*6R>l_X_@f*i>aSw57SlE1@eg}`;17YR8O>P?U|8W+eO_*{KR|e zD{K1VZXP+E6E-v+6Pm6)rpA1ifVz44?r7@CD}Dmg1e15zZq!Kgf`VNqnso4j!j?B9#C!iKB>1_F=F6 z#6l4Vzfq;xP ze{9!T@S29zf%$M)RHc8lI?jpVOcOPQM%SqrP|)(2c; zMyeYbu8Bf~D9B1z^L_q~hnr^}L`<7ymM`^PArTub8VQ^w(*;@HYr%D*e0q`*mg`bQq`sNnVF=8ADW%>EtdVmM=zT3tC)? z8dKP#MZJta#jeB+t5u}EGrtf;pldU4y}&Y!^IzWY01c<|pTQV)_fGty>Z!YhyYRhv z?cJ8@bh#(29HMA#(yGw`$GX6Z9k-UG7>aYOivvXF^%NY{sCsbgf(MICVrrTe{Rm8* zRz)FL;81qoIX5PJepmL!6-b*|h#gj8tVjv7rEP7zDXMOB9=Fh*hY-PU^L`-Nu415H zY)tdU2r(W1HFo*JX0$oMC~xCn-y#9DYG2@>D^Z9CE+jWteSAibTfUu{!CILUF(?+Vd@B8HLnk8<>pjb0wH$(SBa#7MSTycz9f7M)m&LbN4_N9~eT%K`D}> zm*WFw!2&%VTr`IZARvq4!`3~nm=_#=wz$Sh)P{~TiF!@|m`YI$l4In-n2+WO2cHse zEAP9FPLHexT_E2vL*E+#?_64DzwoQrjMRw_SSB^A<8S6iN{MaS#Ao{}S6Y? za4URM4U(FEXBss=@ApGsbX5GptiCW&MedbfYot!ZwQFpnO)}zVywK@CDF_yyvbB) zHm0LROLbY%H{IH-U#-O5P;EYnZZP3Q71^$0oVi}p{evGiU=Kx7YGox4?qg-|jX=453Z{O=WsI6W9EL2vuA@CR{#jo@#?UUv@ zjR!}g9Q5qaKF@;%s~g~J0nv#$qSq5|MMKlfcYC&f8oiZFNlA%{9P?a$slC`Phf4up zy8$46XObyk^cHi2h@~YGgky7t=x}F7`~{twiPdil1vBO5O)_Y^wj9BD zm6PSGZRgL`309wiBSaGHoDD07DtbR=vPm5p)0ThTZw&vbCt8l7#;ML-QXDzOWz4ON znuli4-jQq;pl)&qk&e!RpH&nCYo#oI=c=HPhF#zZfB}Rw>b4z@E zPBl8O!`mZ;1I(Ei`w14pPJU}F~jEncj4U%+{br^`GmF3oIGE`siz6G@M!&9ApPO%7N7=vGcqs+X}*E z$SC>%GLeNP0O7Kb7F!T#aeS2H0rC~#$@D_Xgd;BLrwIx3nbP}t>hspd(ICfy=k={i zuSiBy=2d24L}P@QYuB{2%%4Ac-0KNax>w6xQ3+6?Jl`r&T3j?!ciq^U3q+w==ECi2 zV40?#alYB`Tu~xV)(sz=FN64}BD4&qd*tmgPtsIZcLR;5mfR z--e1*c$|fe3wDgcEs>#SmR1KsPx%vneo3w9a`@5go#zl&qFf!UDfIi9FvrSe9@@yr zJ~;0=L{3FQp?we{gWU|8TAF|&js4MY^j5Cta0sibhI54DZ^QUHd*%c{YxYdBHy!U8r8f9zxjc*_za zFRoKQ(X)8ZS3S*~XjJJ4fF3lZaIjmdebrz|b5EP}{t!hNIu@>z7{GHI=NIrOwY?Es zTv}B5_UuN7@6@-y#nW)ZT8y0HI<4YvMduYNLHLgfbSD%8)nwGn)bWRhD7d4xPtJkG zJ`wWmG!yhJW~@_(aEHnxJxx zAz8!)&DR%Y@!2Ge5PBXKp?eHSVBK2oWSKq` zj`I8Yh5=`mTq)&UApkPn>VIyQ=na>5M%?4j*>h{X5u=3$nU(d669b}GPw!{`st-(S zxjBJfcoWAC5yJGXB%gy;02{Q_Ko}0G+9KQzeAF$_3Xn+;I+e5Af;OL80ILOZHjo;W z%^SlL$KlV;Zf0ena)E+;}L}mek1MWlnOdvdY|Jm|m3K zLif5MDs8<=l(e?yqXX+XAU_*$GczvAZt7e@lDDT9$9WA~7Al0ekRb6WvhXVPj)Ea? zc{jmn4ivM)>jGBD8*34~tU2f&uQG)t%ZV6EPp{as+kB2Q;E&gJD+)l4R@BTC*j9jP zQKHcFWu8Qz@{Q6k;cfvf<*Z{1$61ho8_m8ml9pQ?`9%wRS)J01!*O1sisSOyNmF~g zn3Hb80%g-h9Isc|pW6*JVTf2+j01Yfo~iq*wbP7Dg?V(@&98itsgb;DJ3lIHS*%nF zoabhXL$H}>dFLvd*i7}~vFOLEy2_b{g4rAwaY-UT`WW~`u~_aNQURHM&51T$McryZ zik%lCI4uzF5a+}>?Y!$~2cQWE*!|^K#^v0 z%w>kbtDf|<@5^q-r4qy&AIEKadQQE_)YjoIm8yx6j- zcf_3-CG7&G9F9kJmb&_8mZ^x^S`B77&FJN7$~obUb{I|qhnFJ@p-!Ruw*Z6cLVo2J zYhATVVWgU&bGXOTF1BHt7%}THy$;T~cd67Ul=sUb$z6VNZhcLQ7z05ra;hChBc+_V z;+42-M@Ijcov6A6S2u*u?&hzU`o=)t9`bV>fLdUnSK`9VGPW~ENOy0%1Lp#zz!i;3 zO~!CTR(gvu`Hh7w!x13nHH&~RRLr|NS+x@u>3GU2;{fC}?l5m~mYycd6tibFB*mVi z`6x$!v{ipy50ds~T2>(voHCn%Ov#fqP@FjbF44){em4pP5{g)s^t?8^#N_SEu}bc` z5cU>j3F-qeVxF>=P3+n4eF*c_}~vsIU7_E0PaQa=W}t50;}y#Cl_X4H&8TTQ`KSrC`VLZp$TjlxYhMn;e|EA zJO<0{+IoJ5zMKH+g! zaVYMP1wq7osX|O%#{O>5)5@aa5u|;Nx-~yQxp2X2lWOvJC zK7Am4AVv5EZP3>ryVE5X_2+Ge_Zjz3x_GzK>}1->PWN-We{go#PEO8b%h|z1iG^Kr zN{9O93s)?)F>#HrT1q<2aDgue+X}9LFg88 z>ZokHfi@q^Rszle{!Nm+lp*Gut7J5v`eYv{0|LxS+x;P0$M~?YFI#0yOchB*CUDj# z8PYraHHSeTl2v9m3kMFg{cdQkFi?w0AJ25R998ev9NP9})qciyk|cX#yuZ^r6v338 zLwKa2=PBaA;U+jLG*~{tbn1oq-111SI+UntT>-ItcbP)_WPZ|NPghiNza>g?>+Xd$ za^TxpccVJoy&I>1TCq!IChw_cy6->|gp52K7p$K!dgT}Q*^|_^Ipeq~*Qz9Cj}#0W`R;Ql#aMt+;O!Dk zvvM2lF5-;l!v!)JAaxnBe8k%bvm;-*)1NTjN^7uxy0F--qOgRUU$?wg@?h-Ju9-&x zT`9A1;2;8DR%Wv1sy@1%lb*EJpk}k%eYt+VJw8CT&q7_jWBd*6)Zbkj0w=W)iN9{ai*+%8jFW|2L_L%!N$Hj2pI_e8J!K-m4%Fg8! zlDP~lHb~yVUc=nqudqDc^gONmq*7GjbTVEtH>W8!{5|C;d+k26AC2v4h^y0ysHu=& z^;dsn3gy)4kb3(%|0bvOrM+rO!|?V9S{K+lZ425{ibjf3g+|#`B~v9qGC?*$F##T~ z0Jx;XY-$p5|53^+C@E6|MMe_=GGxHIj$k>}unK)r7x~b^=F&pPjkTm1Sx$!fPYwOu zn44*br6)7)Js&a3?F?UyrHG@e`RVepHNiqDdaFn&ZtK;S*K}2TiF*gfqxN$<^Rugb z-mJ8&$`FH)3Mw6SRMh>KE@KXsR?Mw(Q|lb}4aYFz!qdf#OzB>HD{X5I)+8wp>RX02fg@yZO zKHt5!OL!%(HTZOC@sZ2-5F}C^LIiI)P22dYVxTLI5N2jdL|lB>c)u`Hg5{P)!_<&I0HISGcVUa+j2m<5d;G$zhl@zP7^EXUS@xfyI?A%lWC9W|)NOO0s}i zNI3iucZmZK>pMoLxL5V_@~kOuP%h{K=M7$Uwl&wGTb|$$ddX;P2R7Wjc384v zGk0!qVYn}7OJ}QkYt{BzwzTMna2vVjA7Y-c$20xd$4rWeuBv(7+Qse~REyX~CZ;1D1OvOV4LU`o z{oNuaXdC4$h!)y0S{i0&GgTqDQ~{8InNe|yE+-=YJ0f$jN?4s$a@VryU{gFLo~EkPw7gu8Ajj}5P#=NotJ|Gu0WrdGImj~ixf9t;FSq*wWR-FUxMfK+Yk*k{62#-FrF6R5J!xsCp-81Ee4vF5m-*7p_2xq{t?J>8 z;itplMW;BeVVn6tYNU*kdkl%YzjiV?i+LHGS`1_jrGWo~6X zfE!GV&3m#qg)$YgAcGXN8PBOfu1_ftMeAHq-@CGu@F970Yw{p{%)0jNig#}`?Wure z7ZaR}^JEX%T9Z}EDCtpTb$r;9S2WlsFe=q2qUll)i$+D?-nzW?YE&ax#bCUE5-sEp zQe19O4SVax8CdQ=%0&q!ZLBR!>T{q@!(`;pk8{srEgl!i2R$3;ZON0M3*+m*#1 z*#_9$xSx*B);lMFLu=4d-O`8(LPMcd{_XGUi9rf6&L=rH3Jn7e(szbWnD`sM&GqaMGk0U#qX$#oobp0r4JEe4 zDBA(y#f?s5!?755E{`dcwd1<1Nb!Dsfi=e+Sr*Yd?AoZ0d9hub=AuQd38e><tuLx6B z)OJ^-dgcC;fr55NuUJc0swpXkm0h87X`at=#$T&Hc2m*|A(0rdq&4`Aon&;t{#Sy9 zS!A>opJjYK2mdarwrlLc>8iTei%`mNbSJalP2)n?qkBhoiJXSoNAtoZF1c)@q4754o93O4>KOA zWSsOP?7v%@ITno$=zliTw%7B!%4<3B%5u5A_qgvP#3;#P*^gelL{D-=KhL^P)K6{2 z+3p;L7sA~H)CKG5=@F81fWkgYGyOgt>laV=VBelvV)sqcRW7&lS+#QvzArq^yK;kt zz>x6v{zsPjNCPo7$p_jqt-YC(h+S7jhi?|S4t5@=P}%24U$@mT7|#sq(?lU_Rnb#T z{zEZ`$+efKwf2ceVuOjIsv;}3@Xll3&n~C0Xbu#3ztx81Y-M>y+A_bgiZQ0sm~*lt$@^x3}6wMUOx9cjZ`<<`tmFkC4U*u@xKO@}m zGEYRdoI+?0U(PX|k(!db4xJ;!YB`A+V%36FDoL=LN=Ym(?dP-X=!PkiyE;w{eAFmL z!?isIYf-}ul6J*UB}y`q2hZdD7DjGx#(xy!&-UYT4?NiKzp0b032~wqHE&v6<&>IZ zgN?wD&>hZ8go%+k_HB9e^~+JB5b_<;k|pxf+TC-fs$((pC+YJer^d*z)8yDr>MarT zpQk5ZJ4^RG-L{&B;Fc>S>x2l9`R4DHumon0Rm@T%-|nX za_qgs;WJbM_;vQ+lM{oVqS~jP zp8p&}X0RK*u}?#jePuFCG0S!U(gZFlb(|ksD7+xDRjVmK5&dVG?2dKhVpHOIoqPa>T- z@xHn#J>4;cws59$ALik15@_ah%P1`5rl*&f7+SIi&B*hhbQkn1yLF!x^y*gFHd*j7 z{i}EYbS1)+l#~ixcg;#pQh&V075akFpq~Y}*1vs!xMV&UxM3j}zXqzjr>CbuCI4tw zeSJMBb^!gyjm^z`e0;W18+VtMmXL5xc$tj7{Z71)(?C~J(tVfH@6nEOayosbYsnB?PEHQAOe&WV5)!IsYy3ET3ZFyUoE*70RHdZ60rfhz$<@`u zj~=O@vvF19*XAUqzfC@atX*_HI8o$)R2G8<-4gTRKKKrUa4IL-5!)WgI(UErJ$H`l zlbwe@oMgZYP5u%vGU(a z^Cel#&4fWXXwaM6mfHy#v54}FOixcYcl*=-W*IME=Dexjb(3b9edlZZ6NO2!e-)V9 zCN0VeA0h@;cKylNSvWYNvW4yDuWh?++~td9XJlmLEblD<<2aQeBU`GpFeNjLLo1lKM#81vNAKl;+3qdtbA?W170pn_{D7n zDXl;$ui?x8bGDHETRjb)mwLg&Kpf9x@~o5Pti|Qux*`8fa6yg-C%E9p|3akye)DgS z|8FO*3kRlYvWgZ+2c*cT6AvQ?D!xGxyi3zL@EV^u&~M8QsoYT%-U6liqDLDIpqU)j zmm*C-OdTS^7`~KKHW}tAx+*R}9TY5f@>oa+%as_XAbL0lx<*7QjAdlnK~terFwODd zLZZ5=YDQRspo1Cc*;wZSw`h3XOR1=E0Tuh82#kj^@EF?)Is+yrCqZ|6u?Ab2)x_53 zW`+HdQXVaOu{ZhNy?eo@M?3kPgjcW5QG<@;;^JcW-9dGn_+vp6LOf_y9J>aH_xiHH z+h0$5AigBDYHlf)rrg$R7-XYf#tRM}mmaA)IF#%Rt>o3NEi9;mZfZAu0|Pj_cACel zN|}@IBDPvuF-O^ZEu>8HaUyPNTSz2Q_Osn^E=1Rgo)#sVE*~F8dt1z-S~#d-da=g^ z&Iyg8+A5-x3L4MX^B(UU9^MD-LCeIp!l1!{Ld10^J<%gmG*Y)fU((jLaNNYign~=| zLqzd-)M)tp%8||;4`spVz~{kTJxOI zN}^J(eH~b zIj6S3m>9}I_)T{$bF;E|r`lKFla@iRt*EG|*@7>vhx2s(M?ez{2IJ`Lo>8Wnr6%p= zfmuoi);S(qh(Mfl@=tYibVQ7SqI(rOfBh*04Ahhj(4-d^tJuCWmvg!3Lal0TOJVyE zTv~C{ee@XR2uKa=`P>hbD3Gh0jW{D;8+O^s!eSN2mHsZ+2uj%=;pyRlP$_YHB_A)y zOGoF;PX+XfmtT9_)J?~8Hr#?zbOiYaU{!hY(n9=AjUmh&&?G8h&Cn`&j z#-^5*BDQm_w5JDBhnvK(+t;t>g9dpAx2Dgx))K#9u~_rqucFm!<8Xb<&JqwOQE+v_ z>PS%u7Kx6C_@`oXLF}}myxbY6+|8b?%Y7|M4G_^`r#+N+LOlfBr6T-^#k8}9*3*4$ zTwu}yXGFy76Hj)S_B)g=gaCsY+c}R$vuDuPJiwOVY#h+ESYa1iP0eVmw`|lMF%KQA zfwQyon>Q(LP#Qt|7YT~youHGWv9ZzKC2p&AYakxiqf&jm2khlkS9`ndX-x7(L!cz% zA;w3bed_)D_j9EZDM#S8huyymduH&65-7^3fFR0g^;I0s@lwRu+SyGx&k|W?+-*x% zOnyF5nT13|WP^sClvPs`lTV*NQv$8pdsSrItv3694S1%V`B-b2DCrwAcV*)4h`MX6 zAeWbyfuOY~o9=fbk>M@X9>G;D1l*oIAyUDa+-q7lZrlLqea#K1Ush{ttCc9N*xO%T^vS`&-JtzE zb!^}R)xyNbMyIv$kwa$S!_wsASxg`6)qlWUo?&4MtAJr0_rEFgW;@R``n82K)t+Eh zfCZ}p1+`dlaIcE2zTygUW8-a`vZ&H-=YMFR$ey1NgE+lOVE-Af^! zs@-`DdKR-iui(D*eS30r$Uo_904_TT#`sYaw*UYL1eezfQ<#)g-`(8}<6^{Zyv}=I za4-j;0=zp9hiP%g8-KCX|Fugxb7TLfzOHWR?PJi)3D}H3^Mh8*!@Vtx0U5=oTdFF27sNDU!H@ z>P?mkv`Wc(;|XjP=-w}>ayXbx-bks>Dk)Jtp0T9?BFKsseoWhW0KgNVCPvfQxm+S; zuO%%lt-@|$eqkXCaRjuexE6TaZP2^uvYjTio4OBlM15d!dI_TmPewa}P=?4a2yzJ3CF=l)IHRSwP#mSz6Mnp;k(3>8h}~DtSZ1 z%cgh(%xem|)s|%!vepP>H7}TAXo`ZCU@L8whzb}aUQ5vxqf+ro=6S*s9L^b(eDHS8){`CYepPFlRN88Gk?*};dQgs5rR~gWVIIM^FB|4{tM=&9xIX~> z45ID`dwarB&UZJ0l%)*!c`-^X7Aq7AF9Co7F$4qcr-i)F$B_xWJYHg&R_`_$fHW11 z_S0iFAFp!#x_4*@>Ey(ygg2ahF_lgn5^-t5sMu&v()Hm57FQP#cJMi?Dn(nU|1+n;t=W}v4 z^rca|UOi*pr8ber-Ib!AE7jFDw6+ab+;!B-jn!3E0KL->@K2k%HkJwD%URiAboQkt zuk-Ns+grI@uBYnvkVbNNQh~*&bdc_iWf|n}4Jhk8fHVR(nwgn-k_gVsnB+Owe)SMy z{4M7-vY-J@x|}+yG;V@JvZ@KLk+V?%gaf9w{|xKqkpEpWbKRQSt`->if<)N1Dl%>g z4G%cYE`+HDJS?u?!wvE7-3x1R?`w{YkMCfX_A%`o90Zl;@C8K>!*pa$BkQ@q8j2n8 zz&{>ItWi~=>*qmepFQmpDQm65_!Lq@ zUngd+K2md%@MDy$kZ_LALH!A}EFlx$X(Od#1mbUO(h#phaS~C4KnOoUp){A49sBowLW=rAr|e;gG%l1qh;l~ zD))fK0@NFP3(HW)=kx38PDVvVB`0Gh6?zRAnBTu`?Wx2>9}PILQ0mt6#`f|)$F=xp zfW{IL^LLK@*2s9%O|K0^|5Q#Q3ca@6>1%>%`$VE1Tf%n#yj^q@3N=6!lr#y2&PXH@ zO0%40H;swkufe38mp{^++zo_X@Ll|buv3Iy$}ih8;=W`%HN6W5P=EYOC04U58`uzB zIl+(1%uXpYH;`u82m(-&2%(D!Ct>nkIv!u-QXT9BK5!~iTuBsTHT~GGeu|#>zjpN# qZuf&6`P75|t64y7$i(D1E8FB(@l9VO^i46iH1TvhcC^|x?94wst$=+1 literal 0 HcmV?d00001 diff --git a/src/main.cpp b/src/main.cpp index 53f5600..4276bfc 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 24; // feel free to change the size of array +const int SIZE = 1 << 16; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int a[SIZE], b[SIZE], c[SIZE]; diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index ae94ca6..8c3cde8 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -65,6 +65,6 @@ void printArray(int n, int *a, bool abridged = false) { template void printElapsedTime(T time, std::string note = "") -{ - std::cout << " elapsed time: " << time << "ms " << note << std::endl; +{ + std::cout << " elapsed time: " << time << "ms " << note << std::endl; } \ No newline at end of file diff --git a/stream_compaction/common.h b/stream_compaction/common.h index d893abd..5174556 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -29,7 +29,7 @@ inline int ilog2(int x) { } inline int ilog2ceil(int x) { - return ilog2(x - 1) + 1; + return x == 1 ? 0 : ilog2(x - 1) + 1; } namespace StreamCompaction { From ec6bffa652b7da79cb840f85a5c79cb1b94eb149 Mon Sep 17 00:00:00 2001 From: rms13 Date: Tue, 19 Sep 2017 19:46:18 -0400 Subject: [PATCH 8/8] readme correction --- README.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/README.md b/README.md index 0ce2907..33cdff8 100644 --- a/README.md +++ b/README.md @@ -29,7 +29,7 @@ A same trend can be seen in stream compaction. Both, CPU and GPU methods, take a ###### Optimizing work-efficient GPU scan -Much of the performance gain for work-efficient GPU algorithm is due to the fact that the number of threads at each level reduces to half, so more the number of levels, faster it is compared to naive GPU scan and CPU scan. Mathematically, for array of size N, there are log2N levels. So, there would be N*log2N threads without this optimization. By not creating extra threads, we will only create threads equal to the number of leaves at each level of the binary tree, which is equal to 2log2N+1. Thus the number of threads required to compute the result is much less for very large array sizes. +Much of the performance gain for work-efficient GPU algorithm is due to the fact that the number of threads at each level reduces to half, so more the number of levels, faster it is compared to naive GPU scan and CPU scan. Mathematically, for array of size N, there are log2N levels. So, there would be N*log2N threads without this optimization. By not creating extra threads, we will only create threads equal to the number of nodes at each level of the binary tree, which is equal to 2log2N+1. Thus the number of threads required to compute the result is much less for very large array sizes. ###### Performance Bottlenecks