From 3188dd75e15b8778c60b34c0cf09b6b8f71d4f73 Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Sun, 18 Sep 2016 17:50:23 -0400 Subject: [PATCH 01/10] finished CPU implementation of scan and compact --- stream_compaction/cpu.cu | 30 ++++++++++++++++++++++++------ 1 file changed, 24 insertions(+), 6 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index e600c29..6683342 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -8,8 +8,10 @@ namespace CPU { * CPU scan (prefix sum). */ void scan(int n, int *odata, const int *idata) { - // TODO - printf("TODO\n"); + odata[0] = 0; + for (int i = 1; i < n; ++i) { + odata[i] = idata[i - 1] + odata[i - 1]; + } } /** @@ -18,8 +20,14 @@ void scan(int n, int *odata, const int *idata) { * @returns the number of elements remaining after compaction. */ int compactWithoutScan(int n, int *odata, const int *idata) { - // TODO - return -1; + int j = 0; + for (int i = 0; i < n; ++i) { + if (idata[i] != 0) { + odata[j++] = idata[i]; + } + + } + return (n - j); } /** @@ -28,8 +36,18 @@ int compactWithoutScan(int n, int *odata, const int *idata) { * @returns the number of elements remaining after compaction. */ int compactWithScan(int n, int *odata, const int *idata) { - // TODO - return -1; + //Run the scan on temp array + int *tmp = new int[n]; + int j = 0; + scan(n, tmp, idata); + + for (int i = 0; i < n - 1; ++i) { + if ((tmp[i] != tmp[i + 1])) { + odata[j++] = idata[i]; + } + } + delete[] tmp; + return n - j; } } From bcb7e3f64ddeab87dccfa3d4353fcceabc9152b0 Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Mon, 19 Sep 2016 17:01:07 -0400 Subject: [PATCH 02/10] naive scan --- src/main.cpp | 2 +- stream_compaction/common.h | 2 +- stream_compaction/naive.cu | 62 ++++++++++++++++++++++++++++++++------ 3 files changed, 54 insertions(+), 12 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 675da35..9370139 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -43,7 +43,7 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, c); printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); - //printArray(SIZE, c, true); + //printArray(SIZE, c, false); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); diff --git a/stream_compaction/common.h b/stream_compaction/common.h index 4f52663..73e6935 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -6,7 +6,7 @@ #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) - +#define blockSize 256 /** * Check for CUDA errors; print and exit if there was a problem. */ diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 3d86b60..8aac42f 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -2,19 +2,61 @@ #include #include "common.h" #include "naive.h" +#include + +#define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__) namespace StreamCompaction { -namespace Naive { + namespace Naive { -// TODO: __global__ + // TODO: __global__ -/** - * Performs prefix-sum (aka scan) on idata, storing the result into odata. - */ -void scan(int n, int *odata, const int *idata) { - // TODO - printf("TODO\n"); -} + __global__ void kernReduce(int offset, int n, int *in, int *out) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); -} + if (index >= n) return; + + if (index >= offset) { + out[index] = in[index] + in[index - offset]; + } + else { + out[index] = in[index]; + } + + } + + /** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ + void scan(int n, int *odata, const int *idata) { + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + + int *dev_out; + int *dev_in; + + cudaMalloc((void**)&dev_in, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_in failed!"); + cudaMalloc((void**)&dev_out, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_out failed!"); + + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + + for (int d = 1; d <= ilog2ceil(n); ++d) { + kernReduce << > >((1 << (d - 1)), n, dev_in, dev_out); + std::swap(dev_in, dev_out); + } + std::swap(dev_in, dev_out); + cudaMemcpy(odata + 1, dev_out, sizeof(int) * (n-1), cudaMemcpyDeviceToHost); + + /*for (int i = n-1; i > 0; --i) { + odata[i] = odata[i - 1]; + }*/ + + odata[0] = 0; + //printf("TODO\n"); + cudaFree(dev_out); + cudaFree(dev_in); + } + + } } From 359108824a6a7fb4ac8a8b49834157257b5a2814 Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Fri, 23 Sep 2016 20:04:02 -0400 Subject: [PATCH 03/10] debug on laptop issues --- src/main.cpp | 2 +- src/testing_helpers.hpp | 4 +-- stream_compaction/efficient.cu | 52 ++++++++++++++++++++++++++++++++-- stream_compaction/naive.cu | 5 ++-- 4 files changed, 56 insertions(+), 7 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 9370139..c17762f 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,7 +14,7 @@ #include "testing_helpers.hpp" int main(int argc, char* argv[]) { - const int SIZE = 1 << 8; + const size_t SIZE = 1 << 3; const int NPOT = SIZE - 3; int a[SIZE], b[SIZE], c[SIZE]; diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index f6b572f..ca0d206 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -41,9 +41,9 @@ void zeroArray(int n, int *a) { void genArray(int n, int *a, int maxval) { srand(0); - + int ctr = 0; for (int i = 0; i < n; i++) { - a[i] = rand() % maxval; + a[i] = /*rand() % maxval*/ctr++; } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index b2f739b..c18958c 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -3,17 +3,65 @@ #include "common.h" #include "efficient.h" +#define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__) namespace StreamCompaction { namespace Efficient { // TODO: __global__ + __global__ void kernDownSweep(int offset, int n, int* odata, const int* idata) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); + + if (index >= n) return; + int temp = odata[index + (1 << offset) - 1]; + odata[index + (1 << offset) - 1] = odata[index + (1 << (offset + 1)) - 1]; + odata[index + (1 << (offset + 1)) - 1] += temp; + } + + __global__ void kernReduce(int d, int n, int* odata, const int* idata) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); + + if (index >= n) { + return; + } + //int some = index + (1 << (d + 1)) - 1; + //odata[index] = 32; + odata[index] = index + (1 << (d + 1)) - 1; + //some = some + 1; + //odata[index + (1 << (d + 1)) - 1] = idata[index + (1 << (d)) - 1] + idata[index + (1 << (d + 1)) - 1]; + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - // TODO - printf("TODO\n"); + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + + int *dev_out; + int *dev_in; + + cudaMalloc((void**)&dev_in, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_in failed!"); + cudaMalloc((void**)&dev_out, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_out failed!"); + + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + //up-sweep + for (int d = 0; d < ilog2ceil(n); ++d) { + kernReduce << > >(d, n, dev_out, dev_in); + } + + //down-sweep + /*odata[n - 1] = 0; + for (int d = ilog2ceil(n) - 1; d > 0; --d) { + kernDownSweep << > >(d, n, dev_out, dev_in); + }*/ + cudaMemcpy(odata, dev_out, sizeof(int) * (n), cudaMemcpyDeviceToHost); + for (int j = 0; j < n; ++j) { + printf("%d\n", odata[j]); + } + cudaFree(dev_in); + cudaFree(dev_out); } /** diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 8aac42f..35c2ff5 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -9,12 +9,13 @@ namespace StreamCompaction { namespace Naive { - // TODO: __global__ + // Done: __global__ __global__ void kernReduce(int offset, int n, int *in, int *out) { int index = threadIdx.x + (blockDim.x * blockIdx.x); - if (index >= n) return; + if (index >= n) + return; if (index >= offset) { out[index] = in[index] + in[index - offset]; From 2496e9541659cfe72c0f8da568b9045e6a3a4090 Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Sat, 24 Sep 2016 03:36:03 -0400 Subject: [PATCH 04/10] finished work efficient stream compaction --- src/main.cpp | 4 +- src/testing_helpers.hpp | 4 +- stream_compaction/common.cu | 66 ++++++++++++----------- stream_compaction/efficient.cu | 99 +++++++++++++++++++++++----------- 4 files changed, 108 insertions(+), 65 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index c17762f..f14a97e 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,7 +14,7 @@ #include "testing_helpers.hpp" int main(int argc, char* argv[]) { - const size_t SIZE = 1 << 3; + const size_t SIZE = 1 << 16; const int NPOT = SIZE - 3; int a[SIZE], b[SIZE], c[SIZE]; @@ -55,7 +55,7 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, c); printDesc("work-efficient scan, power-of-two"); StreamCompaction::Efficient::scan(SIZE, c, a); - //printArray(SIZE, c, true); + //printArray(SIZE, c, false); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index ca0d206..f6b572f 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -41,9 +41,9 @@ void zeroArray(int n, int *a) { void genArray(int n, int *a, int maxval) { srand(0); - int ctr = 0; + for (int i = 0; i < n; i++) { - a[i] = /*rand() % maxval*/ctr++; + a[i] = rand() % maxval; } } diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index fe872d4..6cfbf7a 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -1,39 +1,45 @@ #include "common.h" void checkCUDAErrorFn(const char *msg, const char *file, int line) { - cudaError_t err = cudaGetLastError(); - if (cudaSuccess == err) { - return; - } - - fprintf(stderr, "CUDA error"); - if (file) { - fprintf(stderr, " (%s:%d)", file, line); - } - fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); - exit(EXIT_FAILURE); + cudaError_t err = cudaGetLastError(); + if (cudaSuccess == err) { + return; + } + + fprintf(stderr, "CUDA error"); + if (file) { + fprintf(stderr, " (%s:%d)", file, line); + } + fprintf(stderr, ": %s: %s\n", msg, cudaGetErrorString(err)); + exit(EXIT_FAILURE); } namespace StreamCompaction { -namespace Common { - -/** - * Maps an array to an array of 0s and 1s for stream compaction. Elements - * which map to 0 will be removed, and elements which map to 1 will be kept. - */ -__global__ void kernMapToBoolean(int n, int *bools, const int *idata) { - // TODO -} + namespace Common { -/** - * Performs scatter on an array. That is, for each element in idata, - * if bools[idx] == 1, it copies idata[idx] to odata[indices[idx]]. - */ -__global__ void kernScatter(int n, int *odata, - const int *idata, const int *bools, const int *indices) { - // TODO -} + /** + * Maps an array to an array of 0s and 1s for stream compaction. Elements + * which map to 0 will be removed, and elements which map to 1 will be kept. + */ + __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); -} -} + if (index >= n) return; + bools[index] = idata[index] != 0 ? 1 : 0; + } + + /** + * Performs scatter on an array. That is, for each element in idata, + * if bools[idx] == 1, it copies idata[idx] to odata[indices[idx]]. + */ + __global__ void kernScatter(int n, int *odata, + const int *idata, const int *bools, const int *indices) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); + + if (index >= n) return; + if (bools[index] == 1) odata[indices[index]] = idata[index]; + } + + } +} \ No newline at end of file diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index c18958c..49d39db 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -9,59 +9,58 @@ namespace Efficient { // TODO: __global__ - __global__ void kernDownSweep(int offset, int n, int* odata, const int* idata) { + __global__ void kernDownSweep(int d, int n, int* idata) { int index = threadIdx.x + (blockDim.x * blockIdx.x); if (index >= n) return; - int temp = odata[index + (1 << offset) - 1]; - odata[index + (1 << offset) - 1] = odata[index + (1 << (offset + 1)) - 1]; - odata[index + (1 << (offset + 1)) - 1] += temp; + + + if ((index % (1 << (d + 1)) == 0)) { + int temp = idata[index + (1 << d) - 1]; + idata[index + (1 << d) - 1] = idata[index + (1 << (d + 1)) - 1]; + idata[index + (1 << (d + 1)) - 1] += temp; + } } - __global__ void kernReduce(int d, int n, int* odata, const int* idata) { + __global__ void kernReduce(int d, int n, int* idata) { int index = threadIdx.x + (blockDim.x * blockIdx.x); - if (index >= n) { - return; - } - //int some = index + (1 << (d + 1)) - 1; - //odata[index] = 32; - odata[index] = index + (1 << (d + 1)) - 1; - //some = some + 1; - //odata[index + (1 << (d + 1)) - 1] = idata[index + (1 << (d)) - 1] + idata[index + (1 << (d + 1)) - 1]; + if (index >= n) return; + + if (index % (1 << (d + 1)) == 0) + idata[index + (1 << (d + 1)) - 1] += idata[index + (1 << (d)) - 1]; } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + //for non power of 2 + int ilog = ilog2ceil(n); + int off_n = 1 << ilog; + + dim3 fullBlocksPerGrid((off_n + blockSize - 1) / blockSize); - int *dev_out; int *dev_in; - cudaMalloc((void**)&dev_in, n * sizeof(int)); + cudaMalloc((void**)&dev_in, off_n * sizeof(int)); checkCUDAErrorWithLine("cudaMalloc dev_in failed!"); - cudaMalloc((void**)&dev_out, n * sizeof(int)); - checkCUDAErrorWithLine("cudaMalloc dev_out failed!"); - cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + cudaMemcpy(dev_in, idata, sizeof(int) * off_n, cudaMemcpyHostToDevice); //up-sweep - for (int d = 0; d < ilog2ceil(n); ++d) { - kernReduce << > >(d, n, dev_out, dev_in); + for (int d = 0; d < ilog; ++d) { + kernReduce << > >(d, off_n, dev_in); } + //set the last value as zero + cudaMemset(dev_in + (off_n - 1), 0, sizeof(int)); //down-sweep - /*odata[n - 1] = 0; - for (int d = ilog2ceil(n) - 1; d > 0; --d) { - kernDownSweep << > >(d, n, dev_out, dev_in); - }*/ - cudaMemcpy(odata, dev_out, sizeof(int) * (n), cudaMemcpyDeviceToHost); - for (int j = 0; j < n; ++j) { - printf("%d\n", odata[j]); + for (int d = ilog - 1; d >= 0; --d) { + kernDownSweep << > >(d, off_n, dev_in); } + cudaMemcpy(odata, dev_in, sizeof(int) * (n), cudaMemcpyDeviceToHost); + cudaFree(dev_in); - cudaFree(dev_out); } /** @@ -74,8 +73,46 @@ void scan(int n, int *odata, const int *idata) { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - // TODO - return -1; + + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + + int *bools; + int *dev_in; + int *dev_out; + int *indices; + int *tmp = new int[n]; + int *tmp_bools = new int[n]; + cudaMalloc((void**)&dev_in, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_in failed!"); + cudaMalloc((void**)&dev_out, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_out failed!"); + cudaMalloc((void**)&indices, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc indices failed!"); + cudaMalloc((void**)&bools, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc bools failed!"); + + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + //map to boolean + Common::kernMapToBoolean << > > (n, bools, dev_in); + + int j = 0; + //scan + cudaMemcpy(odata, bools, sizeof(int) * (n), cudaMemcpyDeviceToHost); + scan(n, tmp, odata); + cudaMemcpy(indices, tmp, sizeof(int) * n, cudaMemcpyHostToDevice); + cudaMemcpy(tmp_bools, bools, sizeof(int) * (n), cudaMemcpyDeviceToHost); + for (int i = 0; i < n; ++i) { + j = tmp_bools[i] == 1 ? j + 1 : j; + } + //scatter + Common::kernScatter << > > (n, dev_out, dev_in, bools, indices); + cudaMemcpy(odata, dev_out, sizeof(int) * (n), cudaMemcpyDeviceToHost); + + cudaFree(dev_in); + cudaFree(bools); + cudaFree(dev_out); + cudaFree(indices); + return n - j; } } From c3890cf3a839b8ad720f733ab1c7011ade7db548 Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Sat, 24 Sep 2016 16:42:13 -0400 Subject: [PATCH 05/10] finished thrust scan --- src/main.cpp | 2 +- stream_compaction/thrust.cu | 4 +++- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index f14a97e..bc10b29 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -14,7 +14,7 @@ #include "testing_helpers.hpp" int main(int argc, char* argv[]) { - const size_t SIZE = 1 << 16; + const size_t SIZE = 1 << 8; const int NPOT = SIZE - 3; int a[SIZE], b[SIZE], c[SIZE]; diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index d8dbb32..eff1528 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -6,6 +6,8 @@ #include "common.h" #include "thrust.h" +#define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__) + namespace StreamCompaction { namespace Thrust { @@ -15,7 +17,7 @@ namespace Thrust { void scan(int n, int *odata, const int *idata) { // TODO use `thrust::exclusive_scan` // example: for device_vectors dv_in and dv_out: - // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + thrust::exclusive_scan(idata, idata + n, odata); } } From 7303d64309caabc41630c6e4adb2116ccfceb75e Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Mon, 26 Sep 2016 02:03:17 -0400 Subject: [PATCH 06/10] added radix sort for upto 3 bits --- src/main.cpp | 24 +++++++- src/testing_helpers.hpp | 13 ++++ stream_compaction/CMakeLists.txt | 4 +- stream_compaction/sort.cu | 100 +++++++++++++++++++++++++++++++ stream_compaction/sort.h | 7 +++ stream_compaction/thrust.cu | 2 - 6 files changed, 145 insertions(+), 5 deletions(-) create mode 100644 stream_compaction/sort.cu create mode 100644 stream_compaction/sort.h diff --git a/src/main.cpp b/src/main.cpp index bc10b29..259a870 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -11,10 +11,11 @@ #include #include #include +#include #include "testing_helpers.hpp" int main(int argc, char* argv[]) { - const size_t SIZE = 1 << 8; + const size_t SIZE = 1 << 3; const int NPOT = SIZE - 3; int a[SIZE], b[SIZE], c[SIZE]; @@ -118,6 +119,25 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, c); printDesc("work-efficient compact, non-power-of-two"); count = StreamCompaction::Efficient::compact(NPOT, c, a); - //printArray(count, c, true); + printArray(count, c, false); printCmpLenResult(count, expectedNPOT, b, c); + + printf("\n"); + printf("*****************************\n"); + printf("** RADIX SORT TESTS **\n"); + printf("*****************************\n"); + + // SORT tests + + genArray(SIZE, a, 7); // Leave a 0 at the end to test that edge case + a[SIZE - 1] = 0; + printArray(SIZE, a, true); + + //int count, expectedCount, expectedNPOT; + + zeroArray(SIZE, b); + printDesc("radix sort, power-of-two"); + StreamCompaction::Sort::sort(SIZE, b, a); + //printArray(count, b, true); + //printCmpLenResult(count, expectedCount, b, b); } diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index f6b572f..ba40e91 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -47,6 +47,19 @@ void genArray(int n, int *a, int maxval) { } } +void genArraySort(int n, int *a, int maxval) { + if (n == 8) { + a[0] = 4; + a[1] = 7; + a[2] = 2; + a[3] = 6; + a[4] = 3; + a[5] = 5; + a[6] = 1; + a[7] = 0; + } +} + void printArray(int n, int *a, bool abridged = false) { printf(" [ "); for (int i = 0; i < n; i++) { diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index cdbef77..b945583 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -9,9 +9,11 @@ set(SOURCE_FILES "efficient.cu" "thrust.h" "thrust.cu" + "sort.h" + "sort.cu" ) cuda_add_library(stream_compaction ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_52 ) diff --git a/stream_compaction/sort.cu b/stream_compaction/sort.cu new file mode 100644 index 0000000..3d5bc5f --- /dev/null +++ b/stream_compaction/sort.cu @@ -0,0 +1,100 @@ +#include +#include +#include "common.h" +#include "sort.h" +#include +#include +#include "thrust.h" + +#define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__) + +namespace StreamCompaction { + namespace Sort { + + // Done: __global__ + + __global__ void kernComputeTArray(int n, int total_falses, int* f, int *t) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); + + if (index >= n) return; + t[index] = index - f[index] + total_falses; + } + + __global__ void kernComputeEArray(int n, int shift, int *e, int *in) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); + + if (index >= n) return; + e[index] = (in[index] >> shift) & 1 ? 0 : 1; + } + + __global__ void kernScatter(int n, int* e, int *t, int *f, int* dev_out, int* dev_in) { + int index = threadIdx.x + (blockDim.x * blockIdx.x); + + if (index >= n) return; + + dev_out[!e[index] ? t[index] : f[index]] = dev_in[index]; + } + + /** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ + void sort(int n, int *odata, const int *idata) { + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + + int *dev_in; + int *dev_out; + int *e; + int *f; + int *t; + + //int *t_host = new int[n]; + + int *e_host = new int[n]; + + cudaMalloc((void**)&dev_in, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_in failed!"); + cudaMalloc((void**)&dev_out, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc dev_out failed!"); + cudaMalloc((void**)&e, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc e failed!"); + cudaMalloc((void**)&f, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc f failed!"); + cudaMalloc((void**)&t, n * sizeof(int)); + checkCUDAErrorWithLine("cudaMalloc t failed!"); + + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + + //max number allowed is ilog2ceil(n) - 1; for ex: if n == 8, max value any element can have is 7 + for (int lsb = 0; lsb < 3; ++lsb) { + //compute e array + kernComputeEArray << > >(n, lsb, e, dev_in); + //scan e + cudaMemcpy(e_host, e, sizeof(int) * (n), cudaMemcpyDeviceToHost); + int total_falses = e_host[n - 1]; + thrust::exclusive_scan(e_host, e_host + n, e_host); + total_falses += e_host[n - 1]; + cudaMemcpy(f, e_host, sizeof(int) * n, cudaMemcpyHostToDevice); + //compute t array + kernComputeTArray << > >(n, total_falses, f, t); + //cudaMemcpy(t_host, t, sizeof(int) * (n), cudaMemcpyDeviceToHost); + + //scatter + kernScatter << > >(n, e, t, f, dev_out, dev_in); + std::swap(dev_in, dev_out); + } + std::swap(dev_in, dev_out); + cudaMemcpy(odata, dev_out, sizeof(int) * (n), cudaMemcpyDeviceToHost); + for (int i = 0; i < n; ++i) + printf("%d\n", odata[i]); + + cudaFree(dev_in); + cudaFree(dev_out); + cudaFree(e); + cudaFree(t); + cudaFree(f); + + delete[] e_host; + } + + } +} diff --git a/stream_compaction/sort.h b/stream_compaction/sort.h new file mode 100644 index 0000000..26e69ed --- /dev/null +++ b/stream_compaction/sort.h @@ -0,0 +1,7 @@ +#pragma once + +namespace StreamCompaction { + namespace Sort { + void sort(int n, int *odata, const int *idata); + } +} diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index eff1528..1f8742e 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -15,8 +15,6 @@ namespace Thrust { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - // TODO use `thrust::exclusive_scan` - // example: for device_vectors dv_in and dv_out: thrust::exclusive_scan(idata, idata + n, odata); } From 610e792f3cc21d1d620609f55d6b0db62b01da3d Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Tue, 27 Sep 2016 15:39:59 -0400 Subject: [PATCH 07/10] perf analysis; need to update readme with notes --- README.md | 11 +++- images/naivevsefficient.png | Bin 0 -> 13461 bytes images/sortcpuvgpu.png | Bin 0 -> 20209 bytes images/streamcompact.png | Bin 0 -> 19158 bytes src/main.cpp | 105 ++++++++++++++++++++++++++------- stream_compaction/efficient.cu | 55 +++++++++++++---- stream_compaction/efficient.h | 4 +- stream_compaction/naive.cu | 19 +++--- stream_compaction/naive.h | 2 +- stream_compaction/sort.cu | 24 ++++++-- stream_compaction/sort.h | 2 +- stream_compaction/thrust.cu | 5 ++ stream_compaction/thrust.h | 1 + 13 files changed, 177 insertions(+), 51 deletions(-) create mode 100644 images/naivevsefficient.png create mode 100644 images/sortcpuvgpu.png create mode 100644 images/streamcompact.png diff --git a/README.md b/README.md index b71c458..2cb2434 100644 --- a/README.md +++ b/README.md @@ -3,11 +3,16 @@ 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) +* Akshay Shah +* Tested on: Windows 10, i7-5700HQ @ 2.70GHz 16GB, GTX 970M 6GB (Personal Computer) -### (TODO: Your README) +### Stream Compaction Include analysis, etc. (Remember, this is public, so don't put anything here that you don't want to share with the world.) +![](images/naivevsefficient.png) + +![](images/streamcompact.png) + +![](images/sortcpuvgpu.png) diff --git a/images/naivevsefficient.png b/images/naivevsefficient.png new file mode 100644 index 0000000000000000000000000000000000000000..5129a213eb96a3a728cd754b1de848ca542e64f9 GIT binary patch literal 13461 zcmeHtXH-*Z7cLH?s0diW3ZkQkQe@~QM6sYEFpAP4N(2UjARPiBIy4m_DoP6wk)|TV z&|APlQMv?yKu83nmrw!$0_46Y;EaCb{nq_)*S%}qby=>gm~+m1-m>?zpZ)B8_YVVI zq4hi0^YQTsojr5%A|D?=nvd_7gKO7-UsnG)dyJ1y$@1*UL6Bx?0*IKm-X?fOIGV6@iokQD&yl-XR*|&C0l6Sds z@g}W(HYIyEJz3keyD<3nob#gfW%kSlm?G@`7gi3z@UmD*YZa@kIISPU*|S*qD<9vn zQ*&`4U=qg$R`K!eWr1%S!OZy9!oWwMI3M4wM*g2)OceO}MI90FLX8PUG%4CS`{}*a3>_}r?}%H+@qtkLPady5IawC z^tW_&X5_;#qoSpr!tjN07N?BM8N@9O;4U^Kz(VVfOJ z#sAjcerRob6Ngzw@b6Le-@o`7v3U9Aix&=JI9pmn*+NZO?HF4?!tal~hG8Q^*pv9y zTvY2q$qq%k<~edPcZoCA?=jpWi)SZEEIyafA#R=6U2fask#nOX@6zSVm)#Q>x=aN+ z5?I!a*H>=VtM^~y*2k%5e_9>VHsVMta2XUt%<9aO865J6NY?|&6HV1mj)YqGcZEpQ zFJY^t)0zuAX;KT>_6_w77O$@yZ^v>MvpZnpZ}=&+LbqLH=5R}fo|xs$4dLxAf z2up*6RKpkKz zX0In+W#LR#=9r0DxnGE2rxo@%52~H>aRO#T8_AQ>k8S?a;5^#zxwJ4h_$5H3-f8$2 zzwig~;mU;8p`oD>Wii3fbmNTbzOd~I`-&GNLv}w}^zD6OVmD<%w2rru^Iw`7?U>4+ znCWm%J07qmrEE5|?1K~P%bVLC)@4CR@2?MaoBF}YHuN?i8)O5)lo}kRa24h6q>|m!s3YiU_O?UKTPX$UTkqrk>Ji& z3DMc6v$CFu%!2NF3%1t2qq!rVLx_Vq*$0sKg6_;mX>vxg+3h-Q&P#pHEpuihyF;Sq z*7r}3b!uYCky_!TH#?&6$|&RSwotsOgkcnJ@G*J^yCbpA8ZuAkx2 zYDHy{s>Jd7h6!5g9Myb68*7k~J-CX0%~#?s)aHG(AqH_4CM&wGkV8o~A$BHIRf+q` z2tJ%8Ta0Rn2@MxdH2B_UN2l_P)K&s5*|O$q%a}5&%A0_h z`Iw622JO-GOg^9&d5yBkO~}fApfQokg0|Y0Qizf0j%be~Vn$JOt8bg)CkB#mb`;Xg zGsy`>H-Vx%%2I{;-AiN4B{^)&EPwniqhCkj)ZHcpt7Uf>`86%nYcj80)pGs%!I4hj z&JmcoL9eG8;P4x?0#bdSzkR%iv5LIPTsjtRF{#!kynQn&Jhlj4w+E|Nb)njg(}#Cf z@%i$o#m@;TsgQqnDn!CRI)2~t*n7_;7QP5j=cEXJt}in(>ilPz3-!!opy|R})&iP=X=KUf!XVZ}-RXKfcA3~ljik8ejYdd1m-1yh~NLh<4eGi*0 zIc=8?*5mZ+cNW#PS<0j$>rQt`y$Hf{0HQ4DeG2+$TGRGG4Q*k0PxoBCR-N=YH1?OcpDtlmK zHJ2j8lOFf0=t#>=*<)wy3Ds_Ri)>@>LfT(fkZ;G+5sQb5<4q$(Q75{yNHK^Hodi?9 zaF?yx#*(5MNP1I>-qaM`yC=+%&R7~{jHv|Y)6wYzJp_xMt$31he{PZL3F2c$U1UkA z9!6Vl9SlOVb38A~32o%e_TaMN)ub{V5~a4izNQKjicxlT#gCOsvjXCwI(oM)MIWn2!gduAzpy|Gb_ebF5e8c#UkEsFV0k zSCpLp;xIzhcY}QCtW8m0EGjF0v_}Pi>5?>-4(Cq6Tj#FjM@NRgP!_#|Hy>~qd*=pL z(b(!FXWzZ&8)e?fJBQzABtOq01>|ERH!55=4vZ>=S$rmaQ{b;FNDMY*vTotM^xk zAuR3Tp(jLOId3@$79E>3Ypc@UOb8S&N%Yc=^U0NrzZ>yH!rw-Z(F&);aW8g@y)y?6 zDyp;JyxLas6)?q8tXyBkg`>NgVjFHaF_LE|nQbO=GoNZOs-{N7+ifDxkEV8i_)Buk zq-$YbAm;<-=w?puXNE43w10aGL4vcU(wp7s#c(WovB=U6YlIx{0Dvk&1C3+hI4Q;n z@WD-pQy#ImL+2$(o^{ulk$#EwXOv>REchREcu9K|%-H1bK<%_4Jfrm@-RgC?iwv$6 z>w|~Ot?1yd;dZ9ntsIN5C6m=hPC2W#B##V6SG!2QI)LkONQ;kz|8cTVCh*B5txxo5 zgVt2-8mvD`JD+qgW9icxHlIPw{kQ2qCp-_}l+$dQ1zDz7Z zLLl;v;cFx!Sgk^$g7#L@RdVI_#J+XO$6h=0(Sv(H@KfSbSxu4<;F84^Y!hF1%U=7e%~alBQgJFRX(!5VOQ zZHlYF>Onf$x459p1>Un%xoRul3ZMF(#{X|2ktOCi%C*#!fF^}vT z65{i2bVR;<`LgNj#|X;oM2}Z7jp#p}t^{anr;c;w{T&WN%`d3aJxX3w zytG$C7MSPM8A}LPb|=|oHbU(B!^8cQd)tn93X-qB4^HW=iRw?=FRF_ZAHBUAq%69H zZ3QmKof__c#{018zRu{`u8?eQs|fjRfIyn8Mc&sm)b;CGb;m%Sl&BkH3P@zUzB~gU zzWe9PCMoUh?avZh0Z(uYLl{`f^?|!|aKFEH>MR<0b6C@J)=$x5Ftlhy31Kh941rLLnz+0`%E0F<<&<+ zA|Ppi9=;1$&$pKR`&u$YZgnLvIRgt9mHV^mHsyCTIyiGjY^=Szg4_YO&iU9Niq0uY z(Ths0kuNK$?~4h0QFfYH${EvBw6Cu%#mTpf)Ws-D;n#p{v3C=iWTD^ncB)#zrn+ls z!}gDy0XC7pqUzD(jNKKgVSii_&H-5dx zjQmt^w5jh!;98KC!OKz-O|YZZeq1-2f4);^6GXt7fny*%13NK7gU5$0^VQkCy-^7I z4R+EzoO4@1{7h%Dmmy%6hMx#hAO$`)fVrNS@77;G;`_)X)6z`Iru+T;owugLVvM_f zAd;o+4_C9vuc>_4R>oaqwiYfHGfOo)-6{H1E^^B!kkTZBti@Jf4q~G!ow12x#X>jQ zbEE}J>*F=kE+n0{W&3o!3o`TTRbrP~dUbj@KuEMSLNSEHIM#z|vsV_sP2**vn6=+T zh5JN>nCcE=NzW?jJc6ZT%l1{#ni~#Vw*ZoG}`=*H9|)u{|~NAel|#mO9*<>eht6x+x0pNt@E4SPI$_%7bLRNcyxYCBla z6OJ0cbqK<%ymiPsx!3#Sy{wl?P6(@C%MW;2Zl2$=f&8vc9$7SwDqSz%LE~(WB1!~L zV?Dbb_630ie=QlrPSahY_?wlfaY4LvvpGo@#{zjvstE$4i~&9+6ii>3o1#oqNf0_u z6RpWjiVa#3aOTqh;cDT=*3Buw*2q@rIHxhpU>m8&i8HS3!7_yO1RY}C33(duH)`qi zpeZO$v%nbrX~-=V!TY3a4hgq!=44mc;fp;-s|l_7&Wr`SxEp)Wb!2&$#z?g<_aPyp z-D;rX!tXu9qU-ys+Dp{D&%Rm}N~8|~7FB!DFu~H^|JSd3yl4*>@2l~h z+udvnsQQh~&`6}zsoYlUEkGeT#t|fsf7p*}j^Q@x$=l+uj+Jqj%)Y{A+K^$V4gAP0 zz^tNMUYQx$`@iZ8RSpgjNB`*wvVH<$(n~qtb3mD0dVi0misoO>p?OtnvkujDSAk?p^9;5BbP`V8;SYAKo_X8e-O$R)Nnx$cLUnAyf0uT=Yh%G&H>HT5@b`tOg}FH}}*D=TFa0QQetzGDS9@_9%mTWGiYFGQtp&lNu`% zTjK@+&prYJ&P!eL25L%%oAs~Ak+KlpXBMa^xD}-2=;0^j$}u9SWQ^4{Fm=^!l}4$BTjS+6nhB5x5L#WKZ_B z^pWR>(#6+8CvoYoRsWGzfJ`0o8t@cN(0Ry?VU5X>%R%|Un% z&OOlD%(o{@wd4*oz*@461{m>+v^_wK7e;9Y4j^$ql_n_P5q(kKI-a2+jF_%8${Rnm+}nFq<DvRRX&o7-$S8v*w41slXdvC%EXjoXdze1MwMN$L=&5m8)7N`L$u?2GtWUa>`#Thy00C$!@)&Ma z_^u|we}2%R;@-9wd3h>jg@3w@dAXf%o*nOUprqf9I{)0A>p75NS@=5# zd-D2uqIvGiQ*3gn4gN+8ml_rZ8NjVZ6h#?LI_TJ6AFofZ2YO4Eg{i3cl>@2!3$*JR z)5`|`d}W3Dh#|_$A>bPT@K#nrK67j!*I{4}&?AJzlT*fBq1>yuvc)=C#9U#x0t z#t@0Xe2Yilw>FzDpM2Ija+|vw5Ki>z27r@OK*VmAJ3)NGu}199v0a>>=53%gX9Jmu zVfqf%+5daE# zFor{!yZlCTTbJKC<8w3;99?Rte5qQ8JrH;NJ%n{GdKb2pOjYi3J`bT-USU6gF@iDC z+RAm|JYadn)_!5{)ApdD?XC>pgC8s9IPn0RNm&P*DXj<H(#lH-F;Fosu5aB&dz#NTQGp&qP$UxI$X`_o#a#j}sF>eCXiYXg; z7bKYN1GJgiQoeey(|QhVF*KA0D_BG#LFGpKkv4M4D&#h>6 zQjFc=R7-|M)SkxG;7D#YZb6U!#zN1G#Pl@~2pNT;NiL$|&kvN$w4)$k5x?~9rUR_* z^ex?}Z8lT2v`uP)oj~YxSeTvo`-UG-5E}N$7Z9lo0lj1MLJc^h^sOgCg!LQX+!6tu zA5V{;_ZH~$hf@a2641fr0e{bPuiRljqJjlB?l96ua%g;doH`i1%jwK6Ri|1)na@<8 zl3i_dGoZC2dP-GcyP{Ma+rF$?4V@H79U}L&WR}hkxxZT{HR5Sq%k2`uMgta|O@U|! z2;X=(b|r5xJ5=cal5dB3)}Z($KxHADyD0XB*GuD#XGWmxV=iCRT_iHP>>UaUVvivVU%LLrc<1;B%Mv%^5e zZDVo2@wbiet%Q<_=ewnh-pe}F)2KP!Y&3jcm$faJHs1A~603}w zq)d2h*nj%|_?2IvGX{6HHsZB>Tw@jO4Lp_V+WZt%dhYwElcf7=V-&NeZ0f8nRZ>?i zAH#*zL@`E~szXzx^p zUU}?4z~;o?lUXs>e@!}(M>m$A{~M$JbM`+q_vLwhdm;4x<>!|d_HEJMz3{();2qYmi z*-G-9{%itC4p;N1zG(_UR{(a$0z(G0&TnY}%Gdz6vg@z4C9 z_WVE+38u%`C&bhdvDJ6=)q94w&a;^Wj01Y)y%^#xFK9vjb&U0RBW%Ti@qJ#LQ*P{Bd*s-mahQ(l>Mk4?=N-H^Mp?ge1?ZqK`wd8GU~Tk8}NO zK}iYFk|=d!)|v{4<0sSK4T530_SqI=&Mn(}JHR`_2zY;$)fuF}G_>L7q zirh8MmoN#qc89!t8Yt^J075%A2qZ^Gdcxn~-sn`k`4Mz?p1C6We&HYpkB6b|07~Gx zElr>>n~h=F)x|zJXb@-Gzi_AcU1O-)uY5}cWpoSJ(4Cgf?!4&0vh=}12`r<$9lQqvdBnxB&&QRc; z5{RdS7RN;=sJq$anc2Uhh7u8X9KNv{5-2#T#k*EZQ=^_C-0GAEftFWG5aWzhv`Vv> z47et|REs=*MRo->dHf&ZB7CN_-mXB?f>p}!G87#{*7j^3Z2_YpYjs3G%-_p* zpUQKxwR}iX9$iV4Lf<#znK5syl~;jj1c^fc$+taQ0m_s%^8L?71+Kc6brPo_2>I}> zMnya1n(trquUY&by}xlONq69`I9Mw<1ZeNUMnEJjQR>@9hz5sI!CJxPjp$$AM29W! z#qxFD7A*e^A<~trZPDOn2&4Y5KXr_yRC47I>(^TfyRsVrCxmPQnm#m7)1yjuNR{Pt z5ZwS~yu7Gm17_R7uPS!DGfuQ$UJY-JfeArZe|&<~CN`n z53bc#I@aShe%qHMt08K(y{K1EVr#J%6UKSp zn47+wu|yb|#T^@%h}nvFGBz>kPxBgSyGqfA!hVkZztIMWWjeT-p%qGUz9Zkc;qSTx z;L)Sa5bhC;1{h`7nUuka-oc3jL*G}sAhYB(xdC4PW7FD%SvFvwBfg8mKxTFyNDN5@ z(3cAI&fW&Vs#=YJu#klv?#Qy$)98yvdg8gm2(E^e|6K22C}P3J6A#T5bQGlzC{(fc zS!1R@Q=>(o#x1D7%XZr8uIvssqt+KMU!H@?9U4#<64bW?HNK2!?myyFt4|g_>{6Se zeej7=pH;DE3MA0)fjY;AIZ$>AKzHtZN#1)T4i;%u)`aK)^szqEh;_Q;|F%CI$6%Y74d%n)x}htXwf`v%%=R#!W?6vPZiQMnv`2D zjr!4~2V5d|o_LCTsB*z9+b#vfKl?NRaqqqH*rmB`us$r8jWq-X1z%Awm%_6V2W`73 z301Ua52xbMJXSo@Pf8PRz-C~vMT4pM!{~3T(K`4)M1qXj zR!`88CerjGEy3KA$RsQDr4jr!tU%>s+cdlI#c&DWpO8M`FgMu?wg9vy=>uUT?+MzaEYG#u* z+11CzK)To>TzlJnWjA&B(gY5u1+p9BfRs51I@An4fBsB$m>HoC=1>eEn>qM9NgZL} z_^u6T>z?qw;LM274nys^-{52*5jwgho)hOyZ;zW>Juye2st-eo{7v-15lrnSZIVhh ztfbTeSqoKeX1Xo`rjT9djtmXF(+Qqh;D)KvU%w7$3}|=;n(>Zys2+eqEP1XEPvu6* zT12kdu>EyMFPK~TK}krX^v8yn;`41P?%fl2bJX5KcvP8|zR8C%TdCS4RyrPZZ-?9B zaG*iFhArp!-sYATyQ#y=#*3*BiYQX6?XgA+fnNa>cKnRqfU+u98{;%egu9DM>dLj`cUYky(_6l=F-SI!vF~ zdbdQAOq7BJs&-VGHB%XF;ff!{)RKawgfp8#ndsggH<Y8>df&`uv@fz&b`j&iY;cp~tDZP< zzh$SJ-x+04!1kP&r{f#kPGUxo)*q4Kog3EU{@No^S*f)V{YRpW(fiqxOfs8}n(=@R zsPAadzWHidK6BsNX-M^TXnTEi)W0^&N-DfI|N7^b3+}s1x={#@`6*l>OlB=JOo924 zspq?U3hJk@A2pWQ9@Ww6kkc0(im5nOi23}6?w!!{=0K3-)~Hb+(4NAn3+EXXe8l)| z>`2iQ6VzD?dww!hOx^&5g3RncveQ0o!K-%y@s@gFJCgFC?Y?5)<;iH@>7gfZzB=Fn z5vv=H5KEdeAzZMDpv|7U>mA-h>G7^d7J(faOUzfTEJ-VqA9bnSB4to3{ARPnMbOu@ z%eh3smy9*G?&zJJr|&{8HEWaTt)PcP(laT%`2@oO-y8Xh`Lk zs7*#yr8Nw595f3o0nBj~A%jmCI2jx`cgODcFWQnv^$_zHJ&%Thibq2&s(3eo0S<3M z@-rkO!Wu}@X!VK(x?e)ihnEx6@QG5)SYdxEVxe11G#H|YAes1Agy999pIn5p96;;9 z{Xbp=$`zzg{re)MC;jZ=f5x(JPXd|(>8;@7$0z?a$77rv^z=}3;)xJmHG2i zK{#l7HO7`0sC5Oh0BrkoK}$S@FoZ9{U<24xvmpl2at$-5I2h6lqn$Hw0Q zX?4{frBKaErXEV6=%68cVF8qZj$ryV84lxUDDj$~j4CU!2#oil``}sH(;3bwP$|E^VrS($BLz~WOY`#!gN-Ld>I^2hJT^d3a=188 zXP25Z*@l1u0hD6Z!?_FYZ&nE;gI;KOF%#luAaq^<8DYp|6zHE?WVp`{runYy#3nnS zko(1|+%PX%njfB;+8-?L^Y!5xqF1{^TFTs1e}CF=l^K#G>arbwj1Iib0n)?)wQ!jH z`q2UOJyYi(VfmMINtd|{ew%xN;n6Za^M8S0nQ$3Wdt9@dvAhIsm3twlG|+`%VsMV3GL?YC6w>LXUr$ z|6&uUv)r@p2d7l6lTivbP%7TL`|jKzB#(o5HRf4}cPf&ui&jXgFIoFKDO}?)hYQ^L z;XLR={(=cVR04v{a4qL+gke^$hq{XeiC}FMm64eVqKHTJjO+%WkvT9hY<4lfTtNHir zVtg*hdP zf51vb<&%u%2{zzOwiEa*{a&CpUUu3U=zY|_RG#_jx$KQ3Vi$+!24fWMQ;%d8?6XjW zT7T9M#ZaUEGuedCF|Bd>s>m?YSL#P{yeGP4_}<^h?~KftMxC+5f`}}Xzy)P5TYPe) zE4584_oiY!)LLxi{^^uQJr2WmwQ;MbE&dAp^ceLnKdC%5Or5mOe)j=4KoO{RQ9|ki zPXn^_WkXkAIcZ~BaPw;X9ht4D4?_~=2ZR$g>ryv>LcMJgN1*bpg-wRliS94ILlva6 zQDOY`=&{Nq&zQ@LLo~Qox&>X)PSlxrvL3NGWRm^V;moT8 z>1Ag2frt@yWJkHeFVf*Ovs%@%xjLDI?nC%ZS3#I)P$AV;FrFR{XoArGO9L*T^5ln# z-k$zS^MbarajHF`t{%sqzA&H*KRlQP$X*?jAXX1b#u~V@aT2Hro%^!FwOQ9bZvHlh znN#Uyc}HflL-l2Q?mQ7v2MrYzpb(haTU#!HO{+81l_>>u@Dx#kkuJuWt7w03S}cZe zexTWQb`~@cMUM)$K3H3>=4xlD>a^cRMrE<~E{4BnAST69rDe9_AzUy`P$FVGYK%|W z<;o631_Yr`_f?o2DL2{V7g^38+k;)0pt1!(!Gr&4ce(aW5S+!-WBjk7q>z&7ngSanUlva-0m3fUK!rpj776tFb#61(Ye-* zdn-;aTP9@k{`HN^Uip8Cite=9t-rK#-??KN(W?b_UqIcBQA>&umAHD+q<-g(lgXR! zUEJV!{E^y*l{>lLiIh>Cgq0K?wg!EGUTlC1iIk$ zZhoaHymEP?;nz;e(*M<4uU>xm7cX?4>p#8yA8KZ?yZ^_%`JZiu)7}3t#s8bV`R~oB zj@+hv_tgs)R&Kzs1~3UEW;aQTz?w_cW3!lzsTrk(7?Ndu?5U9-&vbecTPZ9ukx=M5 z7Rr7@pnkJ#ST1DpxaW3%U``=*SH@|{oM3Q)mH{fn=?Wg58-9DDiG-1 zeNv`{F75p#V>be^6NBd&%n$uy+>~L+;9q}O1dsuYTyV%;+CSqL6OgA@JSlw!SB6C6U z-Uxy5neh7B*U%kr-!@$M=*yomo2B;NgGC)0TYa}8?eP7H;B}jYwELu*o16D;WB8-L zzC7amjy_$K6Y^D(SX{%R$DwI)TR-INCC*kBq7>`jH_t`%XtG*2A4T5Y{MEI^{Edd$Om*_|#NwU^=U2lA&GL=O@-si(AObwr)(A z`y>+l8qcW0f2(nw=*X!dcN92Vx|6$aHO7aTMeIP&)v1z=c|C`sZZ$YCTO8tvtTAFu zPvU`Ehw2Y#To{p-O8kP$!!Y`Us#$k%Q}<9mma_0{GCnbngga{2{%)@_i_+hognxPT zXMVp?^)0hbXSa!t=Q!$D%=N7EeI}7|c18@+c^iuT60T)XdD7azX-VZ26Ii-K?? zqo0zfGM~J)Dohd+O`98e#_E$~?T#$lF0`V@bGF(@^>gu^*>?x1Y+44}Xz1ZB3)OG$ z?=H=%VhrI~L&8I$hYX);D~`OL16z@!L~N6{{}xm=SwcKmpEThfM)fh z(;antkCLd`w5}nQF!3u7CrcbirR^1dCuWCgb)pnqzX)&mDrj!4jb%SDyZzl~oVy*_ z0|V>_EosJK;?F51BznnQgNCH+Xwh=6b&OYbe2Je8GSp0J%C~D_0TlAZPkRHGg8C%&km{KJM$cfq!X9# z=P2M|@h4rU(Bv*>@-Tt-%ly2xo;Tj)P4dK13K^Ex{T|!&R=2**+uHVAJ7c#cqTCaM z)sa~i)6zey&Ob4E>oCl;nQgDd$$4l7q=%IuE8gd={Y{(t#h~3Y@Mf_B|B2 zvCJRMo<-}Z`mWV-D4D+)Sf#cua8{Q~!W)+@i5<~s&*T-X*!5!!!*&FFi1cIg~)DvbbTa zQoofr`=AFcabJl!*GR1z3T;1$$IjKO8wUF2qp;7NzZ^)sT2;yzkhIiT5xdFkbn4h_ zCF!(@!Xy7>1??-1twt?sd6rSNIsG@yzBHQzQdfIKH``HS@vRqEF6wG#i z>k1g+d@IJ`e0YkErf!YaBU0qL&9%Whv3-yBC-@4xCskTFi=@n#Tzun`u5rlYVbw`1 zQfDtFU+lQO(QVIk&L|7ISFlD1tOM8y8_d^);i`?42WB>4zc~_A(T#fb{Zf`ut{2_V z@`b1@^@tgk>vDG4WU-XAy3y#-46!M7iZB}Ic~Ld)Qq!S_IKQ&zz9<|;crr5bieCQ9 zeL`*9B4nSXK6yOm-?fpQ;F*U+=-7IwFl@-mHA%MFy68@Vgx-jvCvkp1VIT|aA;1o4S_Y-8}L=iHj?F|~Y~ znyS-(?`~>X{}<3&qoQB;@_kNi!XB)Y6w!H4J%8mu8{L7;Nx{=Y8T+0eWXpD z2pO+P$~Sm_tin|OY4Vh8*Zp-L9`^G;0>@Kn#ic`d4348r!Dl#|UN!PL8@`*Y`ATR?CF61OnGX0}& z_1tN6a;EWc|8x0-pR z=@LC`F_&&jeNTwA??_#YAnS$nKlig4D2QZ+*Vvom1HG*>E73t!bNw$(lKfqn6*i1* zdfiHDg8jFPaTm5Y_;qCuY_{HU)Xi4#>28;bjXy74UpEACIM-m;TRfZ9Zoz+a{B8El z?`4Nhs_PaG2Y2N@c>esks|s4EclYBR!@Zb8Y90NQe8>KGDuFrd0@v?3f!nJT$(xmf zxXiZGZqRW~yA)Esv<~^A8CP%-?u-CCK|a|*rcF`Iv()=&>xS3klj5bb6f=pU@psr& zES|oj8Dq@-+AqV=VVwgzJ!hT~z6DhjvMD9$B?}$O*S&uDsi>`aNMTw8;sIK0WptD9 zY$n6j8)CTvMtXM{ZhCppgHmHIt`DfQCYsf(NFp`TC~=$mL;s0>C-X0^w(uTk z(eTOp2i)h9-(i74|HL|Z7~JfH1HPzca&OOM z;aIwjEa{HvBnrF1n7^`sk|+2|>7-U*AGtbhq*mPFA#tQ@`HQl7+8QR7(Me<)yP@4S zh>o-Kc^@~nEw1Qa>&7y8R&RbhF}q`zNSS)TL9p&P4<$J@g}B?sZ@vfJ+?eOQX>i-B zW4cFow%Aq8+zmfa*YsY{Cd8p&Ak?L7<^HSBYx`Yyzp1=1(zk;*qVUbmRT&jNJ`~T) zi^^(A7Syud5BJETNe1H&w%I4$a;M*2&64^Hr9a9!8J57 zns`E6@b)8riAflc$QwQ=nGFp_I@rJECr$@CTn!GaGkha8i3z*Jhx#BJ8d2_SAxUqe z?)2N4n=n~iN)Fy~aC%1t)w$-#{-X=280NuF!NY2Uai=gte%AzlJ{8C^?Yk9JD};8g z*f>U`KaASY7i-G!s0p0PFwELlG`#NX+CEP`B4=$Ck}7O{c`r!KzPh~zcfg-dJhMg_ zm)D$XK#%=8e3T5B%^Q7eLzMeya{KVeM}D-498vF0qoi5r@jISFU>?V&c&uW(=4g0k zD_%U3bO0AdiHco^8kmowgXjK$f{bDvV+s8Jw{PKU_}4N#{Q6IC|LwKETeOZx@YD5g z;ZIhq-5y?}&>t`tHR-t@fm(PRePLoGB{lsl_>b!;Z{BDcV6d)6dG@ApwU7(#q%?X9 zN}9baWiuxeVaSQqk@AV)xD~nqFH>iohibIsNuJIfnHK79$1k_MF|=+<(hXFBRLQGX zhtE%dhGC8xQ0bkyb{b4Nxt2skaHLH}NhpzcqdBDs7TPx9m_^O6J3Uyv)uc< z+qSo_Dt+$y`T5OFSE1xbm0(Pj|=#r z{+3kYaVze3gYFQnQ&Ljijj*JQP-F8BfU6nQ@LY^2H+C7lGMZ#ui)I%mSVzFmMYT2S z7}mQIUpk$oQETzhMD_jhv-ACd_0_A(G zLdm#=cFV>)Ls6|sEFs1LYq&>RrW(FFLRL0~7Q97z-_g(wovaFc+g1-@F0L=lpPqH7fM+ z6{^Fu`7ymj?WnaKy(kezJx@g_VE5{W3XdmC%q#qwQVlW$shw~$0e&>&rXrxs#*suq zbY(CK(~rWKtbWM1#IaetCU~>dZn|$9vP9>CB=UQ`71iVW&5ten`rw4sw{;5N?{A0@ zr;uD*O_M#yKSv&`u03PdnVU9E?Powpd=aG3_t16}Ynq{d`%pq-)mVAvrO}&^Y$bw~ zDYNtY23AyqODUrZ!@)C9(*$0f*-)z*9kI%p}$Fx}EKCv;>VBB1s`#t;{HhwHE?_JbaKe`6u zTH1mO16Wmp*7J|nQR>xEvWfvhxgXaaJ=7^kr;TM~=sJuY<6)bKri%@?+$znrh&p#3`pYTns7d%L$^8V`^s(hZj%nzUXSu{B~S+zW@ES@WnIWEgeD0E5I zO?{Kg5{>>`ObpfFbq{g?F69fct_%seN~$tqT@&509CB`}=kmp=J#)SY5p=H7{Ql8|1fk$q z-|*c%ZFaO(ryPr`P-Vfr!42u#EEKx(V(_vVa4=mkL)^BdW7Cf#onc!Tx z#qu?vu>E*+0&UXfba1Snm_hnRBnTuKZw&VzWl5IIoktS#TnX&pw-+PXjoKQi7t?My zT@Enh7PxIu?XfY_{|T!~*yB*t?Y>7Nup8_bdI@4ZLKPFc=hTlEmsOqG8SIVIb&)Jz z{jiMtw$yocTPLl zPn}2AsnApYf-4cit4X}2hvf8Cd?Z9q`*T7<)dezdnvc&&#MaFgQH*v=LeyF`W6IXQ zI9&f+zHQ8rd-(LHSX;~1G~>MStW7Es5-Z9z@AI}h7Qg+1f9Q>!_HY@ldsS0=^w!U= zrbHdW_^m)!a(&n>5|%8?RC4Sqw=F8QLbA25q2j6JEcv;fR?}Oi_d2PZzMP9Se{|gB zo7GAV(e?fFa_lV%$6`M9KMW9fLG{R7RU8o3Ycr% z#2|p_jI$o^QhoP??ZYSZbGmRk!gq&T_*y~Ebd2;8<*Ur@3_N#%`Hu%mNeS2<*WfK& zNHm}5%xnF?ne+I~489U@1|Balka5$>gi2%XEl1$WS#b)%v2u&yaiUhF+%pNN4MLt z3z@HT_6S2mBTr_Ca6r))-TN$F_WuDB`u@1X2k+&{5DJo^pH#QDE57lNzO&P&^UG%I z*DBA|HRaf}UeMQn&Y^G|U9E>r9BoLY*wLO6kHUoxo5q@`qvR3tAHo4j=Fi)zQjF*KGM~Psz*|vG$4}2#$x9XTM1&L#{glzFfJN zkO6s>^+-c}TAzbDazD`}QVke`%OwzT^iG{Rb?<`<>MjyPQ{ zgkj$(ZUwV#Owx_J!Q5W8-C476w&v)z@z&v@&(HQJ_6MU?U>PLl$66&zsDqKRJ!rF` zhay3rLf34RbWExS>&e>7?>cYTRSJgRaAtp0$l4=69=l79a!iBgeOu-M(>vE@^GepR z)+oy*d?zA7cPsmVUcxxJ7iB|35`G+{fUj-R3lb_Xup4#I%H25ngf(~}>;B3>gsU|Z zeCPat9~CVqVG@q5%s+h-Sl^bd42yloH5+SvW;PcD(|u(5R)HCbd;An4QYvJ?QH*<^ z2q+iwk6#Ime3dv>Zqx13s5OGnO@ZX#+RZ{2`(|Yai#7vCM5_6m`05V~P(YV6?8Fg`lUbL#j))~gd(Ng~6^edhE2ECC9{w(A+2YW@9p9+BG1?!HB; zCIuS`_-HKYJ5unT)aHK}NbWN>EpL+LDOHp2@`T)wt>%T@D8dr2kJ!o>;|mQGXl7ZY zcw{OF4sD$^lky`(9z?w=kOMsJxe&hjOz_BrPi?TbkJM;X)|HPB-0@-IJtz%szeJ{d zC8W>cfj==oQ&KG)bze`FshVtlWZ^T3Db_fy$~0y~KZ?FpXUlqy9EZ$||} zWTq83e6f5XV|{+BqRX(b(E60r)W)#Vxp>pa0w@vr7*LmLIWv>8R!oRQL?JKZh*)T75~yb1{QeNK}qYs%MX>JN7Liaf}--o+0>~P~@fFI40 zfOS)3OnuY?Ka{omw&U>_mH86feMVUd10Rk7f7I@Vt?xTg{_fh> zQKo6!rOR^$zH4{fz>>z&i_$+#FgWgBR99--E8{%-w82@<+Px)20Rk8S5V8p|n3(N| z1GIkcwS*I!{%f)6zgF&!jY;3^eSa~lRf+R;dOI-iVSmIH{S31br0gNR7<=lTc4S*W zLn;=pg_Yf3*Wy{T$+5A!mFJuEtb2;@BtyS`rxjYkDl{wB5kA2~pMol!A?SwJpxKN$ zfYhzGyTQfm(~45m=ix)_AlqdMJ@nGKA=07r1%#N01iNJwbAJe%cYo zqd`7M_66#MVP?P?vK8`y9j;t(_`Bo`0r&t9>w;>gI9caBrFDSBSp$&btaMq zc00hU}1hGedBm$Z1X3=1{D=I6onA z2(p5t$XK1TMjDE8g-GcOY$~=g6C#~um7j49 z5?S>POrO^%&3?OJX!s(%ls1}e->n$665@hW9P{~x!r|3W(y=Gfnll>YiPUOkZ?=YH zNr+m|cEh3B!J4Q@`%>y4QsBvkVxkPVb(tGYqB~Z>x~UpfTk_vnn{nACTncYvA7k#be(FzTapPgcIfWCso$sE|&uwQw zDjzXeOTY{{^8n!_XG^JOf3oe&754}XMzbzeadU0Ioau^L#%}w<4?VvKAyxyZAnl@C zlj*V)B#k%4Q3pn#D5eeYy*L6_5N+po4&l*&z!f{3~8YMaB}Y z_bhqp-Argi7Z1~Kq!oU67gEHP4U*m1A#c|avSQ8F5F#Fl^RSqCkR}z>5+;C7Z9ZAEG(z&#c?|Sj7~F^#pK%2{PM=c1V5-GkcY~*vM;p~Hstl2!*1uP z0V8`?hjuZ-AoeK3G(YwVY&j9Y&ERAPhNxodIbA`e1Mouv@Mlu+wWXQBet9iT1WPuY z!eratMT~C&m<+%X){q?{2&LVM1qHU}!gqc=VNj7|c# zTD~GWoe)CqI_@zA0M>){^A<4@DUs?g5R^UK0%xJxsj088&jFBgjXJe%OD>gq76dKh zQaJX9slfKwa_36L(?fH&GHwv$?7L%BIr-L?Gq&AAAMllBXB3wr&N}Ah8xyjJ~70HYy!qyaIgc0*JeHZu4vm^eNS6|_WvrB)q(KYv?O_q=4 zA1YqA<<~A`HH;pbGteH6g&l)zyXv51b=XJK8<4Ee>#y}9j7|DGpgj@x?`_Et5^DPQRLY@FRct!Nk_xR<1?A5vz_wccFyTJB2H|$(E1xDOS zU35=smHPa(C-97AoGA<_=f$7W_dP_Z`h?~T|Jt6GIrU+bVp_&;*M;5ZR_%?yFwe7i z!8eRiln?96O4DCsE!d(IZy?C%=U!B|T-%edpgufl)@)|EIJ`*RytSXb@f*pDYZg$R zixq$+bHn`_y<_votU46>*OEz{ysU4R8m>lmvC+9xP z>Ysg4?^ZndyE(s)BQ>iIZFN>BY6;T$bM`dN@d%FpL#FoF7LL#7EF@>_ z1J~Q;aBzJej-yNgE6SBTfhVUdw7dQO1eu&(m>gv4MMb2_P_NKN0gR=b*450l@b_6T zyY*WZn@ivrdt3^d^XIjz96RQZXwb7xa!^@hUU$NJvqAi*c1tc5TMzP{gvYjNTnkFg|uY?!XB`x1ojya zJ2V~#$Qs#Yh0o4kvy(E(vqRD4G~k-k;It7htx@DQ@g>Nir})SY9Ob(rFPGPKX^Zc# z56x92!FY*;xk(aQ8LHePJsy)q}C9p7TADv!+&ulIX;w^#nS;v=wl=?le0Z{Y7>{&V7;YzD{PtN z;1Td>|9OUSu5G;lAgza|z6xIb-~}(M#(ipWQ(bY zP4U1~St0_BgNTJ`JWQkS{skbGv6o&ce<&~l;HMe-CQ}5crTWC;?4r0PiEJEgCXDz} z*mLSbwg=gM2RGNTtQSBa3(_o_P{*dw$IP>d5VARd6JE7M)Ppqhx|}qO-d#Nd^_wV| z)?TlN^NY@mUFNq}aoIWse@I-A;=z*X%5AX6zy$t)BReUm)0xjPnq;V93^x610@++! zrLSui!X>WDxVWmX*S}kCILz6zuuym1CQHZ(&$JZ}RKJkk4`Lt4|A?l6p$wmPV#N^6MaD zLnKR$t)&Y&E+5y(zXo=+y@sF5i)sF^fCODH;3&J&@_No9I|8SIAdm=~^fN*xH;*d< zeTSG8-FDHl@Ih4zFMV_{+@gU8of76cDs$Q&(a{Me5^lSsh1yZ7GZff^SV*IVM>5f}NwYLtpnl+%|32)Q+Ke}nev ze`6z;GqiA*ZGRb+!} z7VPQMFYLUNBA+K~)AGY%D1mlsyiIDG9XN~ZWR}OST=b#wR$ew#nLhuUYNMyA&jTn9 zhD2~&Ego!Hap^jS@lrh|IRnhBg_i}4u;w;7#&voGOjA3JyPx@bf@tePZ_ZQT>Y+s|*rbtk} zV@&_sDx^_X?vdcG2p`rn;yiV3?&&GoRkZxEsl-hW>uHb>Nv+wTD1#pd!Z|I|3~f&y zZ_lbn=K9pXaMwD<0SxHlzWD7MEMM()DGy+$$hiE164$W<6_$+)GoEa@KQ22VQeQ2fQEG^ z|BTW`Gw0rrK;hRj4As80?msX!FDvb{uKJGWBy{Ov3JEw#YpBN{@ks5DR10=T&T=+* zpPvGZ{!0Ng@Xe|mUx**LFTpX1$nPJ0Esbj-j|g|i{dl5|xT&o&|5{MubA+?y(3?=W z1Gq@9sqB%b;z1R>KdU`^38AH%{C98LlSO*#OiN`Bf;# zAf}&7{EgO@u`eE6MM=Aqx!aCZ3UN0Pt%U^8{;FV|dn?v_NnpT;?ac2l;%|!|5KtWA z6(Ip{$-rcQ8?t(O>2J^r`y)$0MC~>3^!ng~>h}B%IN?8;Z8ef9R5Ydco;CSrx8gD` zuNBq*0-#_zZc-hjbkiai5SzY+*Muh`>Iug71$)H}-|c|{kv8=Dt?vgtcYu8y;FH#c zEFQ>l*Hkc}A0ht&&%CVoow+mqh@Y$5yf|T^KOw$`A!O5QgfBpP7sCN**_;WJTmmGo z15$!*0BLLk(hXiNF4wIF%1V5I_m53U3WS+Bok3!MRv!-GJQ3xIADzP?fx}0vMl)c_ z;DYNXZ#)Oq?Y$1k)(FU4&w}XIx+~vtqC9t(Me!Pd?f;n2uoMtby#|qn zERaExrmoHU)^#z*N7^icn;|s5{$nQh_Qv{4G^_+@aaMC87y*YRH4sqF*C%Lgi-H_3 zE}+o$v>+oc9{?j#Fa^Mg!0meTqeyhQ?@c@iGcCj6Jm#u&;G=SHHnK%)jtsrz=tnvm zCLt^%5&1{Ozp@msUX1w%JfbS`5i&cCgbWg6P+l+3zWY3g_!WI99Pn~%coDcY#(M2`5d5IcK?q;>6B z4uf_d_vcGQH)T{{5dx=~kki*L=0{a^sAgDV@N4U6oENL=MWAMNf|J`G+wy{#t%#!a z%l+f_?XlPR7Mu5S+BilLOHl-Zz-0dB4sN6~zXtw${0yKY9wTX;59?JUZFC*dD=(*2 zOqkm*C)hCul(XoZLma$Jj$LG52jXwTVsF0G|J`Y^aUP@#|8>eDp?iGHdf*RZR5_BY zFH&<=cAR(6L1bV?S%}~FWdVyZ{OjHV74(Dut;r#-pSzMKKK}f$>=!`t~VIijc z>)JI;M-*pXUvXZ%T5RX+RNry7GIE`sHloi%Fq(1@LLKQ}xG-Mtgrp`2c5@kgVVUsd z%NG{}K)Q~#jGW^VJ>H+H{zDaiVU`C_g(ej6#Iz!0lGS2r?|@Tl0JYGV=TD>xhXi)O zD^HNXe!V*4xmw`Xc)cZHjBr^_xPqfOQb@8{!y=RL_`>fGwvBrJ++<+##LQzNhr!fO zy%-O0?5fZPF0Y2eK(8T&q;tdr>at(B-4D-Xu*0;T=eM<`AN;JOU zs{3NwRk#a9LpTp#7&tGBn2V(^%ym!VFzSL@L<7p_8qU*TC1X zF-nf@)cvF&hY1&JElK^`?$u@S|1^9Kafj~U6AS%q z`BFfPsryEuDx zU?Yv^_eQx+UV@}};7^bjt7?R$UWvEyZU8q*~7LZeC8jNH%?A8EZ2YHUbY<%zb4B zqg$36;q%k_IgVP?_dN-6=dE|qDdeTdECkua2(Q**ZKu!g|y8wR$GJv)#IznYUGB;4tP+gMS76k!atIF_&J}g z^uLMOOtG+0xyxDlnAO~D;5%S|kn6bqy3C&&>vI>KI~bXD`$5nWq+g=*@6lisf$x#% z0bN6e8S?IT*W!=}^qwE_pXWyJ;s-6B+1OUgr@P;loB4H9I;-!6EH~b^djAw>ibC>z z3B)4K2<)mkj`v*O?k=Gf(0`p?&P0pedymPN#r(5h^R)jf{mM_Ha`d-nyU9YeZ(tJt zo{e(O>gEpmzjv@vS)5p*@1=*Qs#-j`(L6;USuC7y3F*NZV>p8?yGy)jQ_9N##&c}* zS|foOk#U2IEcUCu4Mcmrgkd4XLeGAiA%`~eGBRZe?`uD7Um()ELqLuxe+AtN97?=l z)|>!hk$ZQqnhVSj)&=Xhr0}`LbGwU~r&PQHE^FZ}sfe8M5vR#vd%d2n`vCn>-#|Ii z)`e$hd}LXojirb2qk0ESZyrBO?%qF+=m*Yu|FH}(T!o}EmO<&$1_Jv|^;>o=`(=ZC z>GarJ4rtgr#KtYyM$Gg+2eVEx_T^=Sl$h*gmGya6TtO#L5$JL3%%fv$FtEnq&Ie7r z^!f+dV>UJ_UyRZfNT2slZdnLc=xf!d+I{9({O$pUj<|=6aUN$|97lIR3~mI@aTOyZ zZF5dW>Uw;-0^uz}!$Z3x^h`6vA9*ZeCe@&9$QBwZjno4}Z!rdrD$glMkxlf-ycVfe z>pLOqwi>6Vcd9)lRRMElFOFlGD+8lSIo!>u_Ai@;QZv>Q`#h1w2?;Qn$2O9WgjX0b z3J*>X`OU7Rb3VTx9!iKT#cZF`A`M;Uqq}YzIy$oFIR#<6MFpg=?C(R9a#&P@gWqhk z!6S^L3UU{t5?aC})E~@gDm>3cYYYtb;wq?!ZxLO&MvH`U9a{%?8ZN_OE2#IeoC|U) zwMj@7FskDCE`vwJ^ghpn)5uiTjmo;wa3&kQ?j7gMUaRTf_YiL+U(u1efic+2GI3P7 zg1|zriw9OsyeQ?*@nzEdW|Pa21v{2H-}Fvhi{xw6r2kT9aTbhny5^X`q{CO)I-1p( zY|(hJj5oyJbCIp+(?|vs6<%6lDSS}0qoKwB_66U zqVrO8y(GDD7rG?PA~9X&C)ux`;gnM{)^ku>+S!0%^*RE`T8$DXk%ADR6cy~_q^qQi z=EhBLWOk@c$IYj*LIaL z)AHo(-Erv^!`O|z{$8l>JuZv0_+0dj(AbP~93>~MTJmKfl+Z|rh9gULk`KO&Ms*HO zRL?PZDY;EZ9V15nQ8HP|548jjj{b(}Gfdm=UiDq2;kl?d>Is$jeQG`ai9_--?idAS zbZ(urGUUHJUKbN@M&LGfr#eL7>{D=WzSaiz8As|QAAr`sm%bXI5@=1)jkBBDfH7JU zqkk|u0WoRCrMQhfU$q(u#O#g;&$DOK^Uc1`F&_Ah4I~Vy=ga?UISpyK)8?BDvrVm5 z276XZvzEwaC1}+K--Iq;9F-CCG0N$)rp<5al-44g4-rexlW&}j%a0zRBCa~*es}^K zRytBf3+*@YB0RoJgptMur{~koV`I+8 zRpi$YrQ+sYwQ6fOwf-9JQ-LMWl}82_sdjxfC8dEuUGH*e(=QIb@`#>S+h} zaMGJ~u9(-XMZ#Omo{id#C~W7t#yM-Z0@s-ey-{bILFcqr$@78jttvu&&LVR@jZ<-3 z(m3CKy%hIdhPk;S17>t%r`XGcj5gsw{mp%RiJ(wn?JLJGWYGotoI5)=UHl-qJ+M|5 zal}@-1+ud_IlRL;IbBWk4kK*b{Jd6Uu&!t0T(Z{pLQ2_zd3vElY3@OCK`7ree79g- z#$jP(j(C&;s#;uJm(LWI-Bw*b6_HxOi?fel0XL-CNa8$w zovzc+|FmF;FuwcRfXW2SoZ9fhGWx|qF|kJEd>?Sai58;Py3CUFsVJKQ*OB?07U4?g zxCx|tQmZgsPOnOGFWNUwZ(CJAzx+qvZU=DzudlC~b0kV{o-J{mh5&ujt2*kad^hOB ztwA@#5tIEoC8790aAn85ma2o@D-ZmG=`KY|htUC@C`pOJmZti#GTH8&9ISXt%lfpN zcGRBivZ-SEwA*Vu^dcXxg|Y4(_7h7#ol&x2#A2oBr=CZSJ|pyowpiV{a0&ZNtdZ*X z!Y*$&Y;*sZO~FL}sEzohO0ugWrBrjdWcq!xFL6Cl+xD19 z1lIEcoU>E`9l|^6jxl5B?(m#ek8R;(_3zoCw(6!0vN@=_4BfQZcjM;y*~WmP&}$RKbi*eS)Jn+(t1xC&1a@sMkNXan13@QI4=x2MBgff{50?fRK1+pD+zvjO5@&94x<{)VN~_1ph(Zl3b}WwA2mjUgN7!n z;RCqe>Hs)i0kyMzcs%f?BX6>58i4tL5e!`Z4a`qN^#MvqN@EZ+L6xC_GTon5B|Of22eh9Ih^k$_Nmac@RFq?3;>NilSfb3_O9ghI zgQo}_Mn13ugag(;K0QGMrs!dM{3~MOJ*Go~6!zBJ`LAGD4G6KxF5T`}F3@h#e}pB@ ztb}h8lu5+6kQNqxqMPckKY-$yV(*P>l4%s|djD5;_JGjGE+_x4cXi(mzg@1U+j;1` zwBEP#S+-KUUUQ3O+3u=Fo!|XN)A@{Z+M91fUxywY?b?!2lmAU&KiJlxZ2~!EcBlw+S>zE=chSY8nbxmVH2_bbkI2t>xR{;&=0NIuJHEe(BZ%|pu{5P7_+sB zNbcG{434H+XcPcPiV#6r@E$r4qZ09Fa4BU{yxXtMr_ku6=&Gm&)x@Uvp)vv1Qr)mDY&T zDx@#`!8Irq#_J97%lQOfHl;By z>^~yT#L#VdDB(^mkO5^jD>FQtgvTcQuCDydAMCfKX!Fj8D%<6?dElR;L zwhmdt^N|U7v+}}|H($VZI?yV9-n4HX3d6c&`FO#g#97@i-jDZq_-kfs$~~2&exyW5 z=0|4~mU~K%wPxrjyU%?pnK94Tq~6*Hht9PSYNR9I(E`_#Wm%gIzLcTF3T)L{xu^b= zM5%(thUKUDy*3*g&J~uwHkQ8r?UU*1xJWigo>N%M<}}9|ZbsEj^{KynnW{0df<0>* znc!)bU06;EwfYdyBLk;Tv9HF~3mrq`TuNp|q0qTZVtKh6sQao*YWqA6RF+iTCdal8 z^!ZC_i>L*B8WOMHwo~Gv5*&}vUUz)Tw<8mjbkg=aQe3z}(06O#BFkzb_}Qa2#BK0l z!W0Vp(x2CR-vniGA`s^pDLTh$6C|FuSU0EWmrOJ2`%wGMPGx9_uV4=*Rkzz_@wz^; zIJj?5(R_~P8)W8bmh*#NA886-NwzXVB<-aP$d|N z_W1ev)tJdr`+y-yLYJhGNnfY}G_2aJ__?XQW#Gct$<(XVw2)Yzy%H-nmPq6XG9u&T zEIm90uMiF_FW{9rm97!r*{t7hH^xHIH9Un?Gd%aZrrnZKETkHYIWF@?pQ{ZOH*SEM z*36l{7bz|?O*iER@_E-uGlfbdkYcQAw}thLS*4f`zTlvl?``XJ_kx9Ap=~F z46tB3A;%~T&Q3IqH3&buuJ$mK*i<>~QaZL?x_PoaEAi(%sHn9I^|S=W)~jCoQZ-ja zer5(8PN-Fj-;k8ATnA+KtenO(x_+nly@)sospIDYrh_BOmB-@t zLCGPV3VX9A_LTqVzIy@cN|WGAQ|52q z2r4Wm*FwOUWv!Q{8}yrOsm>vmr(NjZ>^J#3JWkYI6*e&hNKOP?e2e6~|3AU-pBMab z?H_Od-)iQc5&l@m+bgPDLFL~)2L3%kM}_+1O|B~(gld9vF-N!fXv6;KX&RA5vw=f| zjmwQwp<*P3oUC!%is0fH_hJy6TLXG@&By?#y1loy5|Ioe(5)5JeDp0DY&esc2y%PV zw>ewk!5BsbPVrGI5aJ?97-E+eq9?2sc>@TG`6f>mD&(HKzS~w5N?XZ zYw2!xm2#R!=%FJcY=d{e*CNMXo1n7X1lU9AsyP(h+$4}m=q&@i_TBBh8nG2P5G-#B zQ}(3XM6@Z5h;n=B98B300z42U$NmNHm(1$3=G16qkAwo}p}~26VWn_Y`u&v~MnfM_ zm+aP*zt#>1k>%iK;)r&|mkE!f*p+jm`5-&UUr4fsSOr4~TS^5XPS0z*h0$g)${p4BcPo5y~*RGMk0c02f*hVt4n$SETw_m!Wz2Yr|4 z_ZJlsTGit|mgdq|zVhgX#dUquG6RwkMRf4gR?|U98;l&=0bGv1CdUdLTh{*;$E0;l zFfYXJEp2%04f)SdZCpIss5A22IIPCZZ94DZ-uhGMpSTWfv zTwDbajxABC@Q9!W#FsG%Ds0AV6BH(pQ+`OzkKutJelC-BH7)^omwty~I1pFUJ@gS# zD5=*L@U;t-yoW z!R^rER+;+1gKBcM;Lq!J>$I*)QqPZqqcdy2*3(dVrT!OxJ{8L-p2-#;b^H+HGFH-% zHdzY3rm)uH8}3^TDdbxJMUOT3i`m2KF%A?nZ*FHRMR1_XWR2T$-jcgU9L`oV_KtW28W=25?DG{V$Fi<27x|I|O=^nP| z&>@LTujj^}yy@A%&DJ-*}k_8+t5n)|+T#ktP4u6d!Zsz^h{M1?}3Xs%wl zd>w_Na7LkiIelP1w9H{IT|%LhqOV@QsOi)_(|0iJ=4kktg_!V|^kmc5y2Y*zES9~p zvg+roQ#20KKl03fk(=cIMTE~YMI-t6QBfnmlM3%n)3qoSUA$~-euarWI^|$oC#TRo z4myg12fV8T78**E$ASv1Z)z08CCm}Gq?(-5T$|EMqHY9@V+Xd}=jTtsEK)1y#rDGg zX_X2Kkb-2&i|X<|EXhg+kedX zuRs5x_y4D9fAU!9wQEXjRh8h~dA_^A)-u0Knor){j@`8rpD$f%d3{QmNy>Vzw`elq zjcV#d(`Kh}VB!r`vBbd~Qo-)(K#La|6V~rCot`RP)ev1WuvIotm#-uty0O(*fZJ5t zHQyB79UDmCGb~QnnlCR`iOq@+{4~>*)B63>lh)Ynn%D%~2LJ9|@$a4#**I5fG~dkE zrV}#!i($qAVoA+A^X0o9o5(mKTRS%R4HNODdcjsU(jc2jkFII-w~^kw4r9HoYb=W; zsV8>*S0-D}xbDR6%$SlCDSsW&P-gX5A3TkIxxs+eu_;|t=z8`zpGnHa-Zxg5FAUA2 zqh3}X6d~p~YA@aMz*^760Iw2tA?u`Ee@8iKg+b;uCQ*C%<=FWU4^tnXUO&JuTCjaci%<67l&`K+JWcA7*)p9hp3YFN z3*l8}b6?hQA#IGe?yj@#a`dbfHMwUb$zc*>B8Qr??Zz8(jR-#WU3Ww8t~tU9&0Ojq z*lrjQdRk=TKCN%QK^!$-8f%!e*7)?cucg zxu-`}&71-_DWtX+KIfdzEtGZJoJ?jetIRMg$@VRl*=dz&U_FZF`238fjt}F@A7xS> z_DV`ol8W+*D*cTscEp z54R5dUGU*lo8^GmcvR~I>FuTGCTnVuBD^ImvCdz-=_bo}*LGvd9FkR|vp*K!tqakT zS$cZ#Xzi5X^OAYIL=Ud!RH)Eq%mF8mn_D@srHfUjYh~+0zBFs{=lr&a0|#eU!rG;8 zD9g+zZmZs#T+_AKJtQ$Tpu;8OiZ-;@@RhLu+okzFr%76(1m!KJZY#iymg*_6u~|Hs zn&=z0&A+{~wbnGZyH&osH6upN#uoU3I^SUXF1~2ozkCQ?Ze4%ztlPCrOc~QJ_wY1~ z=i&%W8;Ho)!XfA+dY8LY$zv8>eVhHqD)jURupR?4yW_j1&BjAvge;3X0higl4krJP zBl=HY>3%@@Wd0v_y6FH`kH&iwruV+g)!20M+?a|?r+q9(bv@J<0 zk=a{A<+c(^;Y-w$g|}?x7_rWy!MymnN0h8#Ql}LfY9|GsAK=g0eQ_qntZ+UhKA={? zK!208e*DE2mGar)>B9NI-Sxnp6#-XfqS8imTg!Y-JBh`itvu_x$11Dy$Sadwojsl} z73uZ!?%kZ~FLRBje%j=*nlss=%&!7QP4xr%`G$Ka*2CFh&xu5b!6wp76BFOD_`Sdp=gkP^)$K%Vl1s1h8o#qC%$3$tC;)wUPrUZO^3kvxe zxO7V#tRke{T|RpmbB$Zu$k>$*=p5!4)9^nom^5gboZjtXtrr=~fXz#m#X!S#Fs3Fm7d6slC}%-nB%e+UEA4U<2L%jjmB2F z3l#Up1agf^39qp#|JbSg?Tqf&-Abj6uDTEBsgzyz944mhVZ&|PIxHUH?xqir&dXiD zWQP_|4&}45$4QIdC?&k3@-kw*?#EhL!K@)oANXQ}gU}S)bmm2K@HRf|5~I=O#60Sb z?I8iaIVC#u)y4I&*f3vtvmyL05lgC*E-&A981!0}+5RkZXN|zArXQf;krsbBmlwTC zSj$@xM2CdQcj+u{oxKU#wb-QTh6`$nb{cvy8EzV5zTD~d^r2c)Uh>-ZOiq(Y;6?=+ z|FW-9QI)E0+h|{H;XIDiv>Q$$^6#GKUh+F%D^+s|Gicg>BYcpwU9X|Q82Ca)UGqjY z_o(A_Qz|hUc%)e44GZ_qTwD z20Z9zIgI9FL*5)r(UBKS1=Mm*2{@i*T|?e7oPXK|eGl!~CY*8N_6%YE?NHM}5dYx1@%@C-47m*~PGA2kdwJ#p|2+Gy z@5AM9_iAoGX9|+pShH__6`Ds_I@Wk|x+}-@v#}AO5EM_^o4R)1ZCpY{;ZryF$f$CA z6(_-8v_H%+fncHVUO(t60dMvtx{iOs7=LZCH+<+O<=Uwcf|o_T`}V92F0%2$u-|Dr zqLP@xhTL)N&XjhJG0Rfb(CXXgtldd=1^ltjSUC*Hu|DPulka%d&>5MHtD6g-X>*8O z>$*)8@MFT?^hth-w5s?LTJ;yVdOnrCPuZD)Mzel^Bp?3!5Y z5?1fbbY}75Go8v@=ZcaqN!zg=F^;i*&(5zvTFUMjij9_B$n7vza#m~DDGen`yiXi9 zs4f2-ur8`q+k{~?Fvwf(N4Grx9W0?lW^x-VgB=MzGt^|=5d43QziB4*h2T@^dTE(_+Q z*PV4O3hTK9@9Sw_kH;$$>{?TJ7OUQb`7(|RDq+D4BHJ0n?t1(Oqb3Tnu~HT4*~QAvsr2zGKBCd{6VuOVPWvLH4^lYo-HiJ*{do(Wlj7 zdabu0Ok@ufxq!qyA-wub)>YMdekt?-jkTEsO|H3&M>|G5Eq%qIm;`7}xHllqsX@3S zx5op8lC3sov+VrHPQxZ$+U$A*!NsYXY*8=fFOJH!$pV$9Q1wY-k^3KHzMhGTT#- z9BEz5m2Oy)KZwA1-i8>dlIGzE?he2C2622kbRE5wu&yW ziIIDYL-GyM3X4n{BA)e3cVw!?D+XE@z{B=Yv@+Z0m=9c-D;#b?JV{;|5~VOpCEZ9? z_0PBNE57TyDj%#fXdKbnvSa1Eqhver_BoaF3|nS7$)$M6hnp$*&fL|-EB%*frGloW zrm{U;1x|{(Z!XU`AlnmX{G}~AOyQ{S0)cHbzV92q$D*(OOjebNRi+-*y)3;=jzG!R zeC69jcPi&ZyhoxL?^i14+{mRS8C)R8!*s7w&UwEON(#zD^`z`0cAG6)Ch`ZX7)nYAk`7>;AOOXRYF8&kLeIK+l z?b|dI6ORj;+$nqKx;*jfko5X@A&d4jg+OL z4^}ugc!jLTyU+KRt#OS6Z*I!0f1+)9LMt7%;t?>rHL&ZEFYY#I@#WZC&KwZ2z9yj#mg>E6wzmgW2=vIxs zZ60t*bkky|1yZGxqPzfC5OM;-Wq*Y%z3V41(wh_UJZ1#WPW;o->sXnJTZ7x>DaRy8 z>%$PR`OtG`XAW-+t)|}m5aCX{xYAo>$823SyxP{!jl&!4{c-?jwUdUh5eJ-OU8;Rw zf4+VRgJ-e-dQtcN@Ja`Ez>KP~pWLs)n8RFe*L?{*pTWp)nE3m+UTgE;&f^2<$}0{# zexXwoaw#(lq~mZC82@N9fcMACVAdcCXCMO6L$n(UH%$WN&aBy?>yL3B3m5SUJz$x2 zU+FDpI!{zVK|vFZbHF$28V`Rl3}R=exE+hOXr6 zj*yq45b!?h;nvQ$R&~Pe+UHru#O=O$kB<9W4Wa{<*E70WPCiEV;p^LXmWx3*RIk1{ zt)hgG8X_GtKq4qfxvthLG4ve4Kw$AY-FpLUo+-BtpKkVVG|VDAD$(K(W3e*bHuU_Z zN7B)&A|o`ck{6gS`h8plFjAty?afGVCDxqDX*!%4q{GzbLN+#=E`?3>*(;7dJIn z#jJY^+sa(m(s<>TJsZ(_U6C)*7td?$_vh7W=Q$A`DQbhsKdc;dU*UJ(Dh3bS(|Tq8 zGNCO$0hc1D#i3{JM^9g?t`!@$BQsN57mGqPsSu`3E(V#$0q1|9smM+xHbQ>*hTZ~A z)!V4D#Y*NwkCDgI^-Inr*2JnzKH7DBSmahbtda5Vt>fFBBH~CBjtx4g&r_?bxc>tZAM;{L3)2D=OJkj|$l%9x?i;)R{z2l`#96ir ztfMCus0xO@qOnNuzB2WVo8@7fhIysm@k#vFoU)Y5@)b%pxBC&o7FzM+C0lbPIaRB3 zeG+{hTeIe1(s}W#+?l1XpV2V48-4f7WP|Y2zT)7SgeQl@Q@T31N&VLcrt8B6bGF1C z=c14^60+$pRcLzHC^nt|dV8~)0_A!5ZUo(ek-k{fV#?FFg5)^idu=JpgU{Ja_MBCc z$i#0>((Hy8!ZS2Rs3+e<*1tEt3K-*(ejN3A@i05Df}wBe4Vh|C^4VV zSK=t@Ts9`H_2O}ePrj4->c+Bg( zv(oV)AXTpv=`+3u!m z?><6Qa?i{$ZECcpWjyz*KNx4`X6|G2+_jQkk|@+Yy|lRSP;-Tv)8FXe!5KQm6R&_Y z(;`k+x)lzIxw^W{UI(ac&2LyN)DSLsCK5-+#@*<7?*3R11K9?j^9`kja?6$Rt8)Di zONIl9q)VJus2KsTDZw)``c5dv30u$QN6>$YjMtv_X=RU%0=DX1i~V#vZ`n$_zV64a z(3fc6G`D%z0p&;$tJ)X~Frnt_3xmK`#9e&SVv$+a&dv@zpvPpFX>5bSjM2A-p?!u! z?u4zQV53s53OqpXu_2mX?=I|JOc4UTxMEo1pki2T-Ca3)GuM-Cqjw_&&d@3egY^+| zTKJyrW+UKJaR2s06&`vRm$N3971lELYt`d?-{xHKGQrbzbd4(Lpc~rRo-vgn&z!gUB)d)M#|Ht!fB{h`dTxbT;`s(Mk3)z~ zYkoZ7K(~U?h_uYSwZ0^V+ z|LK!&{{a5n4&U5AfdEKZFCX<#IOk3ZYjxqI343#|KE!UT+4KuwsW zlEirXVz(d;YgG~L3qm+o`_9|X6LT8Hxutfum*|}(hdU1 zb_%}UP?HJ{xG?0s`~ka@N{Gt`Ra~y-2u!HsFL#TT>jyfk#d8lu3wCz}n=EY6o&a#o zWX}@NV^EOOnmAU4v*MY-p(NJUw zMhMe6AQ#&Jshy4$WZgo@`uIR;aTc&jldyr;5PGn){C@>dz7D<<66A07d##e%6Z81W z(^Ovl{F%ip0V;%81|$Y#qZ*Lb77hb*eI?9z(kEJvm#l6JpJt>8t2^^J8AI3EJb#RF zTZ#st;HPDrk)@kcA9CobN`K2+$lBRlQMM^t(dsUBwqqcpmeF_rgj!-CUgVhwCh8(4 zV)X-K-Uo?~(%*Rtn-|wLN3XiztX%1cdag5yZ0>dglH^K}={0-Azh+r@2OQqTHsQ>d>e}g<`lBj@bpuketLfM&7i-(MK@w8H6?L zaBq{=B{dq)A|^lN-8DO+>GUNGm_~;g^TFV;v8=q)W7(s=FGq$}!H)F}t$i`;0&Lzk zLOP)H`8lmkp%j>o)1wgt4ri-_&P(G>HicymMr-Zs1*(L412k#}cIF&u3I^Pk8sH%> zb`5X^acT7z-b;UPI5DEr7E$UYXZzK+hmcOMbddi>(zP0`FK%121v$Z@p**^X?+k%( zse5Sca@Pn~a@1>x(u4YOyE~Q{`nmTffW6T=AO7LP2aD_lrCdyB&r4A-d6s(xD!L({ zn^j2v)#w<%7TUjQ0nQ>>u<=5C^V$GJSGsKFoow;`Uh^CZCyXzKR>&f&|2Xia+7OfU z-cW9Kf7y`9gl$QoBj!Da=9@EWjdIuhvgMt8>5AGfi;VA`=o&_fu4pseMY6)_t)dxPbwL>Ef@mo8zM#d^}SC$w%%LnWI4E| zSY5i@qMXyj7$543=C|$?yGFC%t${w*AdojT;H>dTumi^O*8e0SIQSNZ_sh33Ec14Ggm*k1TQo?+8>8 z&LvNEbdQqlvbsi6I9Sepr^EeD!!(o^kw=a+Fd48dm@Gf!S-&i_#Fs0h9lgELgp1)T zU8p$ZJR4$g|I@y%ts{t2u@HCwhT7U-&grYb+l)>SlxPNijDc1aN55JrKVs^$>EH5! zg%ERAXZ-&eQpHKCOl!cp4n5#8lRh!aW#-Q$eAaz^@VTz1PfH{4R^nE@g<0&JqcLv8 z8X+J9avx5&=R)MV{NJ#{ARnMa zLCr28-UA6Ftd|WI@k2X5e8|KT^q$N)&rL7wxO8?Kf)9OH2iV;v>rB18)UK*Ske;H} z^iq}CS!Tc%6kTr^>hJ(oZ*3jpyB55OG0x?Gd`lUF+gVBF?m7nho+_Jr#oH6zQ!5rg z>53*&Df10G_@xT3b%L1{<=W-Jv?Z!b;x8OyD0g_8mk&|Zz%=Anb9yND|Le+dRbZ3C zG8Le`^J?S+qKkKUNZC$V3XU8}p8PLhRhRcIfHe0Us~+zhB{oVf)#rjCc#RcRI^B_p z9k|WTk4_9HY#FKE=IK9)h$O=wUZe93Y;R&gzY90D2ndwXtKfO}c~^7bbGqyx?<#{xKHTBxvJjdoO zdJApIJ=%bI_Y{rKG`{E@DTDK1l61PSi!Wd#=ofx0e)b7e(~~OqFWjO&XC!3R{7(Mv zO3cvCzfN5EI2Diw8m)et49w_Oe5C&B}A?N2DcEcr~eCX$yR^R1*D^b-E0Ys=q>ly z*|Y~yN@l`sqWL0<@@tiFE4{f#oKjkwhRppJJEP?H%@IkJ;~z<=Ki2+V}fK5l>5 zVHm0kiPF+}O>N;BHgS711SEW( zg&j`1{rOoQ@NaPs>84KO@})<3ndY4!l-$L_`=EE`-+zD9h3_;v28rr^i_U>^+-@#_ z?q@xKSQ_UnGGPOzO``uUbuX{DCcA&?F3hzYW12sNIogUQ|nXic_~Fw^vBrQoRwnr9EVd9hmpHZlibNK+6$S~1Za-vgB9 z*X*B&loI;o8tDTBV7|FL*_9fqk$pTd(iY2CD>@Afn!CVBbfACFY#tlR472>n*Hhb3 z)*UAai+U}GS)>*y3j30TRw4WZY>_W@p&ldnz|NG3yAQcHV;^gGRP5SxRU500SpK80r)XZuc@^@S%`_r@1hYuUqyF~av<}(cf1emHQ}p=kk#~E{D% zO@(CXUac?G|8vPA!wE2PyO+kWjx7Q|ZleyN94^$MWJ~_q=^nNH|6+WJ0Opx#=cTbk zy!%$tr9)!B(XfaM0d%kLqC~AK!(0Ja zRU)U5`L8cgrH%vcqO-k4g1kU&ASEm2;H4&k-?{=}O)j!oW_p2e{^N4Uz_ki#Nn$qr zlBUN|W^aBV57gZV*hMW?C@Ha@Obks5#kkaj??)*e`IBaXX=R^xS)LGryn+3~C(2|* zI0TF<6M&z*`_=}cF=XLa{JOuf=&&sq3DL?{b9Z1n8dWUq9m0DMtVwR{hEnk^X{!~$TZF?*=im&5tpDFswAAM-5cLT6j>zFM zguzK5The{DRVQE2Vp2Vh=z*nvwvDh ztN*|#?-oHT?}Y+304v3#_uI(q5hGPrvRV2>rdgY}Py9G$I;d{@X|_{qA);%!scjhS zvN{u9Y(Lp@e>Y0ZF3K5|a-DnJ4Q&6ZNYG>Y7oMVUkjs!hNmPv%7w#PZw=9Gsxq+!k zz4lY~7}l_pR03J2kJ<9#*6m^hY3tLaLIB zmyySHet-ba^D3SZ2b}@GIf&fT{$SKU7>Zv`JiTVqlh3OXc8(oMGavhb$rLxN+je3xX2vqX5Y2ubG$C9{;C?teri zj13fV&*b_MXY!xc^3Ed@gqu-m`S9&-s9{LI^LVt$CFCq;@7CRbCA@ z>Jcw#A@$047@Zi6fT~q3prZ?vbxOW2qEd=pqCc27Mv7)8som#dpD!AZvCvfQ0(6;X z>&Dlu_Qjq2hC-0}7w6d=)gHQ{0k*J@bh_tCTo38p%dX0V+pw(h#9uXV`4d#A27& z`Sp~0LGV<@&L|Y!ovD*Y{Y>tD)%r7VN3~g2bN@?_;7&|+SH)PsD@?s$AiJUnM)3;% zN`SJ>3t1jr_{9zqrk?rU2r-MerF-W_SiqTd;o{M~)cmFDE8OP!!Xss*&FM)W?u;v7 zcdb`lSzRXHYHDkTuC)4kO8$|{^~?`LP!YlcA!FAmzB)NV7~)29W#W)M&Ss!Os3VJa zUKKHN{Tt)4v3=Om(V@=Ebx$S)7-<@e8!@@MW`NVxQzxE;Hhh+d^8oyHPiQ3X*tSVi zO!ORNXqJOB-A+wMQlf$~&0Zbe!=E2)cY@Y4;*hfv#hy4K1*Y%O${zlRSPJ!F_Dsl6 zH}aq~Q}zw4-`R@PGd-m3Sj6@ z%mPN2i=lkBrPO%|$-qQgp*SyjjfM+u>d7CKt+rqvA)F`4Y}BZ=6yR6V?V*$+35j^2 zwl@(;!6TupIdF|&6zt=Mi1lWnP^{izRLz}lTotekg~5s7d=u;VN_bcR#UM`|9t{W6 z7ikX}((rwr=MyE7QF8p5Ma6Sa>ejnBOa{_Ifl93)PSy${5H{{?snFD$=92IyL8cIi z0TA~g)|*R7A9V7qlsJvRTJ={{b%N7}u9(ltRRJ9Dp+p?U=9G4M1k5*Pxq_NM0kD-s zE9zq`XVc!p#BRf<#X!I%A3S(qBr_O&zv3V)DEH`SxG=q;X`^JjiJG& z?YR|mt2auN*MA*xFD8-Zccq!LB*5UeB05ec1Krp9zzzmj-l0Wxx_UhkZQfFKHu2}0 zv)&!La**w5soKED+n-w?w-vsxI4l%#D{UZ^ZSQX+BdL49kRL-sR7nYvt>0B>^z0pn zHx2q`FIM+p56NsVj8AoMMw#=|4 z1yEPVZV`W2A&@m114Q`DG-IE{%*W60IOi*3J(HHCn?!tOvjhOiMufaZ(C^IjJZ^Vu z5-h6{KG?q zlh96TrK-nmi@oH>g*1|i0LWRgGoVr))(S?uky|;DkX&@Y%}y)(xced`@j~%S(ermk zBT(woL8k5J$Pa0Pl&Bm%qk0VjjOxdI+Z#(cV-p&wP{hJy%H|nyvk?Cbw&MuLL$7qY zcUeVmXyq`c#HLc>5rvNIW%K>$cU-Wti!$zX=Av&$N1TG0fRjFz$TM$s?&s)1?b7yZ;7Fp2=aY{wWB4hicYPd#mmYuS>ETR5VlF=pfpc`ojA)Bi9Xb$cat(D?HvEVsu;IOY3obu|i_*v-*QVX@!lT86591#&%#CJ*SRJr7t}?;mg_X z9cGdI^HnasrYO6&kD?+jOfp4MAjUoxWR*fp_(6-X2qm&J(m!T|=*l-mXA_#U)`;vQ zbVEvj#ZKvUHLyKjJ5Ka&zk2%c07#XO`rBzGTZR6zn1_`*Q-C1~HR*<&{Uu zD#ys(#TQE6V9q^@wwt9oTGoDm+;$O$h`ImsoZf(VN)Cc_0`d(*EVicakaWj>yL&TF z_4(+LA;mA}VFr~pY6U+SBsl(j>qc6hM&Wh5bjMBAp9A02Nk>~)>2Q_%(9?F@vqbW? z>SY~U?cFI4ey*Ayb~$8>+VwVL#y(_ee6VVs28$n$MuX^d*6>+cW}x?zH=HXb0Ia+Z zNFCBs9~Gs&$)W2py5B%hk(V)qvJ91~RN@=M-Wp5~CFQhw(j%`?uj$d2DMz+Fps9j!hg1U~Dg~89Q$4#SNWt0+Yk)j5n znVwwp8%z=omVi%lT96DpZ$a;C1nvWosGfoJV(}E&X~jI&q`n@vt6JK+Iba~bz;W&2< zF0@zx**Y5Mq|_co9c-$VUQs!ATJIM*)^BEj`pf~`PJ-jDAFBu7&KIc|It@%EQg!%w z7IICdPzgvX+uU6ulaB4yJxy%7uUW;N>>L_)5#F(3voez<^lUE%HyNIrm7#s<_pQD+ zVtV#6N%zEC&Yr?hhk^5-l8VX<5Vq2>f5 zf0UQ31*u&x*)O#~Ajf2F)0GW(g3zNONcOm}3@q&F(Xg#-wm1$@$+t(-5xvr|fP$IU zl&JGc!j!bEQ0@39mkDWsSa5rNP+qxOAc|{qT0xv66$sL=Bi~XHYyoL}lUbUR-@(RR z(mhv4m@O_YP74xL7$QftZCwhwU0H4D(U1ueNG~epS2?D~0lJZA0*U`gC}pZ) z3K_WnnPw6xggJv&fK~F+BH_twVC9$>br%*AcMV$= z#Y`4Yp4A92rweAf(fJ9i6c{eg+UDzGCYq!2VoCCv_ZBGSezqQ^5F|*=5+7W9A|f2I z4AVZeii|VLuP|z~Z+qxL{e&?~$I^bwdy<{>0sY4-_QCWxv9`R}LJtK@!0`{?OG3~y zOe&>8%gBOiZ=`wZ8U!4Vd~ofP>PH|rP$`E6&Z2o6eCh<^lcO!1|7>kAY=U{_3f~Ce z>OAOx3`zfMEj(jrD`sNQ#2-7ie!$@|3$h}417tU)edtv`uIgcvtY=;I6nZ&4ys@G|IWy71myZ0s>}@SnZHRXPD` z*Y45Luq@h7hmb8kkB)!P+?r8%LeQnVC?7q8BCo%}Z~*4NoQte<&yOMY)x*YizB_S4 z4c3CZNmW!C2$PO+s1?{Jfh^-}H8D8?GhBtt`F%3X4VlJ1?;?AC|9g4?$2-=^4?_QE zC&bHPKn4w6p?+cfElSu)S$XdV^*Wam$;}$bsPMA-9wK>wjfc{YSrDd${z1CTa6-~( z7`1gh;Tfnmrr@*_I9u?YS*1?EbopUA|Bs#oB1-zT|LFFj z(Z4U|siY^>;rIu>bJltRVtFOg?yx{;^P!WP^~u~YrVB9&!&rOZSNV;* zdy9uGcE&Qc=*WFf+{!wjBrT5aSR=&S5!`3qU3Fg0_7!trQ2+V8tN?tuyfI<0_zRbp z{?nofEP825tI;+cOHVazMLq=2S~?C(x(TAi#*n;*&v^4_yq&GPjkF~ei z8_eh5c(btLcGrc)^z_GXI=U+l7)m?0KustdUe#Ian_s&p3_PUR;_VnNy?^|sU(es_ z_G|tV5f6z--cDTWHo=hVi4e>k&qzg|!n-s|6)MClxme3X{u-mNPk&6S=Xl)v@5}gt zrO^Z@kV8*mbBMwaA)PAa1k6bGA!fKB?^U727M0PDh81&Vj>XcOkGM4HU0yeY zA`+3p>e`wr%dto$`Dv^U5hKmme32vkmN#$xo>U*sPYL~~m-*<^=O(@?yckJ!WeT2x zMYL54<1AWJsr2a!+xhMOXnR<$Fm&^an-~bzlkTtW3PWD$RDO-YlLv*dhhU0vBF75O znkjSGesj88`SU4;jv*)~TRMzy8n<%IVDAoQe1|J?D3JSe3wlsXQbTrgA?H+^0gN*h@A3Ir#jh$`lsz-CxXh&k@#`1AZy8d=wnD-4YmH-0pqy7&HacDnLz zLE`1}^GivWk03gxr%||!Q$F>2E1Ud3XS3E08l-*o;@)&M1GB56H z)wb*^{k8|$u3s+SLb`=X{JfC<>&V4P5v%T?$^e$gzjJE|w9+3J;32rjneazm)vB*p zxF^PAb5h8ec!v^h-anN*+g$yMqc!_(9r)tjLIFb{xDNe`{MI|Tb(#c4C!)RmMRul4 zB9>tk6o{+_r@%!5^NtKdC>@*ckM8?1MjRxF?|#ZSHv+dkZ^g#sXlG)`dCV?L7W|`^ ziYh)28sNeJ^p$1&{DBGzmXf0RCofb6FAKRWoBc7RpSAlffUUCW$|hTrO$*5PC#aCD zQQF-K$c5xQn`mZ~VoD%S>F0g@zc%GR{)L=P4kQB7_`f3s^$2wgkm8LrTo!(|t!UyLOB?e%alNz~J z;sTkGR=5a)pr96VjV+{1p6wuaj*Fpot(8W!_v^&g&=3*A<_D`XwMMqPNI@*XwXm1__KUBU=B-L0#rJ+m!|zE2D*$7hg>voXmFV#b#P1Eg$RH)HN2HOIL1=Qy*%ollf(34@=X zOLX4#=-1m?yRs3TbG?xZHa@9fXuv zngi}Khk0~?Q0LkM6)RSF*RAp0Hm6e60}Y}h@ZjQwcklP>ZN^AO;eaavDV*ti75qk{ z|HsnC#VqPPz}93V$@V0mj=8s7AXjf*^ZHa*2UJR_B7F6mDpAKr*_K_|_HYs4&X$p{ zMk(x1)hFIzJBCDMTFC|3)sEQ%PT1Yej@DSxWGoD6-|tg?<^&WS{!R21ons5yNmAGQ z$J2s#K<4fmLyEoY?fEHO<&htxQZ`(!9mo zFQ+Ol-ae>XZ|MC>s&*}k-nU5ZOpMdJCkCg-(-Y@kJZPxc$h2*KZ;-Ry6!}uSXx!2U zHxH*h%t^KU089#ztWj--+UdBcab2lc(eWMIgav!lsO$9<|1Hbhk6JHL8IFYp)&$h#< zKgq8o5^pa!7w^`jyGz<(W~g&Uu8zblOnt}{?__ZU*@)rv3TQ_!&hscR4hZ9$?S@2gc*Se&U1N(?5Ou+1+0q zlmGP}28sM9Sf&5!nB4YH9YdSjS&YkM;&LLt`}@IFgwV!7W==$=slto;22!RA7eynG zoV*Bb7x{?=Am6nes>b0*G_F8GYU#-@FsSU^unV?E#lry-*_n|08%8c-K$UM2^n%<} zMM$p%2oz29_2Dj;g!{?|EN{v15)|!8?+`=Fmw)!LUaO6Ny)W$}WQXi4kDCf@6$W!_ zCuO?NTD*t2cM|kY6|VDd$0CJqmRR_a0w${llL6)cwid0TX zIY3M|377J(ps{YYQ!RG!a!RJX=%ws*j`e>J__mhe;n;LrM9`4=Qca_OE z_mBX8b^s|0&d@s&I1ka}X8#&K)GGrrR|#3hRY=-eNYVs;a6#An4p$({xm&WuNX_N% z^n-#JS1>BNH4~Y*ooRAe0EE+_%Ld4=Cor@l$*38Dx={WmIZMBLP&8I$rZLVYzL`fKol~f%f*%1{%k z<*#G7m*HXwl2)cE;5bn^@g2G10k_&>k-KgR18i^k6%HXdGH&Lw?lAHjCC_>Ulvbcd zLfF~B?^!w|f}7#OKl+sKAeadABhRR!0Yx+fFd(dC7m4wr#8{ZIl8a@~T`2M>7%YO) zhlpU^F(Vw1mDwNLyX^>941;2!R4aXO3vzgUmKdW#xCGG(MFC9iAqU>rO1`^e-(TwN z;f-J~PFgFj4nKN9<1cAn5tjAiQ6p)6ztz0o@Hp_TymPSs2z=n%2jc z#`%gyUMp-FXJELRR={HC4nDLKa0wl7gv^7q*UrrP-rdwQ1TM+ZuUVqK1FImhmOY@E==l(8` zs_HFv1dB1>vJs=i=CD}wTM>ngciVz!4mXV?rXlC}YsZjSu2$}8a9AWm{B`FGn^3R* zb;cNOC~opxAl!?fvqr8(1Xg>G}49nlnp|_Tr3)-O|dxn^4hg(GBUb%%f&?N#&vk0f%PG$x*@yglckJHT^dr@H{&xr|n0iHDFO>(1tj>3k z%dUEE=?CPqQt!IiNGk5l;#{jI82qktDkrF_a>5{OLD#9|e(oF;vR;#LnBkl(GUCH7 zUcm<4s$8q`d7I^9nyXo@eA5_jBoq)> #include #include "testing_helpers.hpp" +#include +#include + +#define ITER 1 int main(int argc, char* argv[]) { - const size_t SIZE = 1 << 3; + const size_t SIZE = 1 << 26; const int NPOT = SIZE - 3; - int a[SIZE], b[SIZE], c[SIZE]; + //int a[SIZE], b[SIZE], c[SIZE]; + + int *a = new int[SIZE]; + int *b = new int[SIZE]; + int *c = new int[SIZE]; + float time = 0.f, totalTime = 0.f; // Scan tests printf("\n"); @@ -32,42 +41,64 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, b); printDesc("cpu scan, power-of-two"); - StreamCompaction::CPU::scan(SIZE, b, a); - printArray(SIZE, b, true); + auto begin = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < ITER; ++i){ + StreamCompaction::CPU::scan(SIZE, b, a); + } + auto end = std::chrono::high_resolution_clock::now(); + auto duration = std::chrono::duration_cast(end - begin).count(); + std::cout << (float)duration << " ms total, average : " << (float)duration / ITER << " ms." << std::endl; + //printArray(SIZE, b, true); zeroArray(SIZE, c); printDesc("cpu scan, non-power-of-two"); StreamCompaction::CPU::scan(NPOT, c, a); - printArray(NPOT, b, true); + //printArray(NPOT, b, true); printCmpResult(NPOT, b, c); + totalTime = 0.f; zeroArray(SIZE, c); printDesc("naive scan, power-of-two"); - StreamCompaction::Naive::scan(SIZE, c, a); //printArray(SIZE, c, false); + for (int i = 0; i < ITER; ++i) { + StreamCompaction::Naive::scan(SIZE, c, a, time); + totalTime += time; + } + std::cout << "total time to run naive scan, power-of-two: " << totalTime << " in ms, aver: " << totalTime / ITER << std::endl; printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("naive scan, non-power-of-two"); - StreamCompaction::Naive::scan(NPOT, c, a); + StreamCompaction::Naive::scan(NPOT, c, a, time); //printArray(SIZE, c, true); printCmpResult(NPOT, b, c); + totalTime = 0.f; zeroArray(SIZE, c); printDesc("work-efficient scan, power-of-two"); - StreamCompaction::Efficient::scan(SIZE, c, a); + for (int i = 0; i < ITER; ++i) { + StreamCompaction::Efficient::scan(SIZE, c, a, time); + totalTime += time; + } + std::cout << "total time to run efficient scan, power-of-two: " << totalTime << " in ms, aver: " << totalTime / ITER << std::endl; //printArray(SIZE, c, false); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, non-power-of-two"); - StreamCompaction::Efficient::scan(NPOT, c, a); + StreamCompaction::Efficient::scan(NPOT, c, a, time); //printArray(NPOT, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("thrust scan, power-of-two"); - StreamCompaction::Thrust::scan(SIZE, c, a); + begin = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < ITER; ++i){ + StreamCompaction::Thrust::scan(SIZE, c, a); + } + end = std::chrono::high_resolution_clock::now(); + duration = std::chrono::duration_cast(end - begin).count(); + std::cout << (float)duration << " ms total thrust, average : " << (float)duration / ITER << " ms." << std::endl; //printArray(SIZE, c, true); printCmpResult(SIZE, b, c); @@ -92,7 +123,13 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, b); printDesc("cpu compact without scan, power-of-two"); - count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); + begin = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < ITER; ++i){ + count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); + } + end = std::chrono::high_resolution_clock::now(); + duration = std::chrono::duration_cast(end - begin).count(); + std::cout << (float)duration << " ms total cpu w/o scan, average : " << (float)duration / ITER << " ms." << std::endl; expectedCount = count; printArray(count, b, true); printCmpLenResult(count, expectedCount, b, b); @@ -106,20 +143,31 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, c); printDesc("cpu compact with scan"); - count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); + begin = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < ITER; ++i){ + count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); + } + end = std::chrono::high_resolution_clock::now(); + duration = std::chrono::duration_cast(end - begin).count(); + std::cout << (float)duration << " ms total cpu w scan, average : " << (float)duration / ITER << " ms." << std::endl; printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); + totalTime = 0.f; zeroArray(SIZE, c); printDesc("work-efficient compact, power-of-two"); - count = StreamCompaction::Efficient::compact(SIZE, c, a); + for (int i = 0; i < ITER; ++i) { + count = StreamCompaction::Efficient::compact(SIZE, c, a, time); + totalTime += time; + } + std::cout << "total time to run efficient compact, power-of-two: " << totalTime << " in ms, aver: " << totalTime / ITER << std::endl; //printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); printDesc("work-efficient compact, non-power-of-two"); - count = StreamCompaction::Efficient::compact(NPOT, c, a); - printArray(count, c, false); + count = StreamCompaction::Efficient::compact(NPOT, c, a, time); + printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); printf("\n"); @@ -133,11 +181,28 @@ int main(int argc, char* argv[]) { a[SIZE - 1] = 0; printArray(SIZE, a, true); - //int count, expectedCount, expectedNPOT; - + totalTime = 0.f; zeroArray(SIZE, b); printDesc("radix sort, power-of-two"); - StreamCompaction::Sort::sort(SIZE, b, a); - //printArray(count, b, true); - //printCmpLenResult(count, expectedCount, b, b); + for (int i = 0; i < ITER; ++i) { + StreamCompaction::Sort::sort(SIZE, b, a, time); + totalTime += time; + } + std::cout << "total time to run radix, power-of-two: " << totalTime << " in ms, aver: " << totalTime / ITER << std::endl; + printArray(SIZE, b, true); + + + zeroArray(SIZE, b); + printDesc("thrust sort, power-of-two"); + begin = std::chrono::high_resolution_clock::now(); + for (int i = 0; i < ITER; ++i){ + StreamCompaction::Thrust::sort(SIZE, a); + } + end = std::chrono::high_resolution_clock::now(); + float fduration = std::chrono::duration_cast(end - begin).count(); + std::cout << (float)fduration << " ms total thrust sort, average : " << (float)fduration / ITER << " ms." << std::endl; + + delete[] a; + delete[] b; + delete[] c; } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 49d39db..42fc04b 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -14,11 +14,12 @@ namespace Efficient { if (index >= n) return; - - if ((index % (1 << (d + 1)) == 0)) { + int off_n = 1 << (d + 1); + int off = (1 << d); + if ((index % (off_n) == 0)) { int temp = idata[index + (1 << d) - 1]; - idata[index + (1 << d) - 1] = idata[index + (1 << (d + 1)) - 1]; - idata[index + (1 << (d + 1)) - 1] += temp; + idata[index + off - 1] = idata[index + off_n - 1]; + idata[index + off_n - 1] += temp; } } @@ -26,15 +27,19 @@ namespace Efficient { int index = threadIdx.x + (blockDim.x * blockIdx.x); if (index >= n) return; - - if (index % (1 << (d + 1)) == 0) - idata[index + (1 << (d + 1)) - 1] += idata[index + (1 << (d)) - 1]; + int off_n = 1 << (d + 1); + int off = (1 << d); + if (index % off_n == 0) + idata[index + off_n - 1] += idata[index + off - 1]; } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ -void scan(int n, int *odata, const int *idata) { +void scan(int n, int *odata, const int *idata, float& time) { + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); //for non power of 2 int ilog = ilog2ceil(n); int off_n = 1 << ilog; @@ -47,10 +52,13 @@ void scan(int n, int *odata, const int *idata) { checkCUDAErrorWithLine("cudaMalloc dev_in failed!"); cudaMemcpy(dev_in, idata, sizeof(int) * off_n, cudaMemcpyHostToDevice); + float milliseconds = 0; + cudaEventRecord(start); //up-sweep for (int d = 0; d < ilog; ++d) { kernReduce << > >(d, off_n, dev_in); } + //set the last value as zero cudaMemset(dev_in + (off_n - 1), 0, sizeof(int)); @@ -58,6 +66,10 @@ void scan(int n, int *odata, const int *idata) { for (int d = ilog - 1; d >= 0; --d) { kernDownSweep << > >(d, off_n, dev_in); } + cudaEventRecord(stop); + cudaEventSynchronize(stop); + cudaEventElapsedTime(&milliseconds, start, stop); + time = milliseconds; cudaMemcpy(odata, dev_in, sizeof(int) * (n), cudaMemcpyDeviceToHost); cudaFree(dev_in); @@ -72,8 +84,10 @@ void scan(int n, int *odata, const int *idata) { * @param idata The array of elements to compact. * @returns The number of elements remaining after compaction. */ -int compact(int n, int *odata, const int *idata) { - +int compact(int n, int *odata, const int *idata, float& time) { + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); int *bools; @@ -92,26 +106,43 @@ int compact(int n, int *odata, const int *idata) { checkCUDAErrorWithLine("cudaMalloc bools failed!"); cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + float milliseconds = 0, totalTime = 0.f; + cudaEventRecord(start); //map to boolean Common::kernMapToBoolean << > > (n, bools, dev_in); - + cudaEventRecord(stop); + cudaEventSynchronize(stop); + cudaEventElapsedTime(&milliseconds, start, stop); + totalTime += milliseconds; int j = 0; //scan cudaMemcpy(odata, bools, sizeof(int) * (n), cudaMemcpyDeviceToHost); - scan(n, tmp, odata); + float time_scan; + scan(n, tmp, odata, time_scan); + totalTime += time_scan; cudaMemcpy(indices, tmp, sizeof(int) * n, cudaMemcpyHostToDevice); cudaMemcpy(tmp_bools, bools, sizeof(int) * (n), cudaMemcpyDeviceToHost); for (int i = 0; i < n; ++i) { j = tmp_bools[i] == 1 ? j + 1 : j; } + cudaEventRecord(start); //scatter Common::kernScatter << > > (n, dev_out, dev_in, bools, indices); + cudaEventRecord(stop); + cudaEventSynchronize(stop); + cudaEventElapsedTime(&milliseconds, start, stop); + totalTime += milliseconds; + + time = totalTime; cudaMemcpy(odata, dev_out, sizeof(int) * (n), cudaMemcpyDeviceToHost); cudaFree(dev_in); cudaFree(bools); cudaFree(dev_out); cudaFree(indices); + + delete[] tmp; + delete[] tmp_bools; return n - j; } diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h index 395ba10..ca99f23 100644 --- a/stream_compaction/efficient.h +++ b/stream_compaction/efficient.h @@ -2,8 +2,8 @@ namespace StreamCompaction { namespace Efficient { - void scan(int n, int *odata, const int *idata); + void scan(int n, int *odata, const int *idata, float& time); - int compact(int n, int *odata, const int *idata); + int compact(int n, int *odata, const int *idata, float& time); } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 35c2ff5..34f8268 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -29,7 +29,10 @@ namespace StreamCompaction { /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ - void scan(int n, int *odata, const int *idata) { + void scan(int n, int *odata, const int *idata, float& time) { + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); int *dev_out; @@ -41,20 +44,22 @@ namespace StreamCompaction { checkCUDAErrorWithLine("cudaMalloc dev_out failed!"); cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); - + + float milliseconds = 0; + cudaEventRecord(start); for (int d = 1; d <= ilog2ceil(n); ++d) { kernReduce << > >((1 << (d - 1)), n, dev_in, dev_out); std::swap(dev_in, dev_out); } + cudaEventRecord(stop); std::swap(dev_in, dev_out); cudaMemcpy(odata + 1, dev_out, sizeof(int) * (n-1), cudaMemcpyDeviceToHost); - /*for (int i = n-1; i > 0; --i) { - odata[i] = odata[i - 1]; - }*/ - odata[0] = 0; - //printf("TODO\n"); + cudaEventSynchronize(stop); + cudaEventElapsedTime(&milliseconds, start, stop); + time = milliseconds; + //printf("%f\n", milliseconds); cudaFree(dev_out); cudaFree(dev_in); } diff --git a/stream_compaction/naive.h b/stream_compaction/naive.h index 21152d6..04344a2 100644 --- a/stream_compaction/naive.h +++ b/stream_compaction/naive.h @@ -2,6 +2,6 @@ namespace StreamCompaction { namespace Naive { - void scan(int n, int *odata, const int *idata); + void scan(int n, int *odata, const int *idata, float& time); } } diff --git a/stream_compaction/sort.cu b/stream_compaction/sort.cu index 3d5bc5f..eaff952 100644 --- a/stream_compaction/sort.cu +++ b/stream_compaction/sort.cu @@ -38,7 +38,10 @@ namespace StreamCompaction { /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ - void sort(int n, int *odata, const int *idata) { + void sort(int n, int *odata, const int *idata, float& time) { + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); int *dev_in; @@ -63,29 +66,40 @@ namespace StreamCompaction { checkCUDAErrorWithLine("cudaMalloc t failed!"); cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); - + + float milliseconds = 0, totalTime = 0.f; + //max number allowed is ilog2ceil(n) - 1; for ex: if n == 8, max value any element can have is 7 for (int lsb = 0; lsb < 3; ++lsb) { + cudaEventRecord(start); //compute e array kernComputeEArray << > >(n, lsb, e, dev_in); + cudaEventRecord(stop); + cudaEventSynchronize(stop); + cudaEventElapsedTime(&milliseconds, start, stop); + totalTime += milliseconds; //scan e cudaMemcpy(e_host, e, sizeof(int) * (n), cudaMemcpyDeviceToHost); int total_falses = e_host[n - 1]; thrust::exclusive_scan(e_host, e_host + n, e_host); total_falses += e_host[n - 1]; cudaMemcpy(f, e_host, sizeof(int) * n, cudaMemcpyHostToDevice); + cudaEventRecord(start); //compute t array kernComputeTArray << > >(n, total_falses, f, t); - //cudaMemcpy(t_host, t, sizeof(int) * (n), cudaMemcpyDeviceToHost); //scatter kernScatter << > >(n, e, t, f, dev_out, dev_in); + cudaEventRecord(stop); + cudaEventSynchronize(stop); + cudaEventElapsedTime(&milliseconds, start, stop); + totalTime += milliseconds; std::swap(dev_in, dev_out); } std::swap(dev_in, dev_out); cudaMemcpy(odata, dev_out, sizeof(int) * (n), cudaMemcpyDeviceToHost); - for (int i = 0; i < n; ++i) - printf("%d\n", odata[i]); + + time = totalTime; cudaFree(dev_in); cudaFree(dev_out); diff --git a/stream_compaction/sort.h b/stream_compaction/sort.h index 26e69ed..4152c2d 100644 --- a/stream_compaction/sort.h +++ b/stream_compaction/sort.h @@ -2,6 +2,6 @@ namespace StreamCompaction { namespace Sort { - void sort(int n, int *odata, const int *idata); + void sort(int n, int *odata, const int *idata, float& time); } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1f8742e..e7f8e7f 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -3,6 +3,7 @@ #include #include #include +#include #include "common.h" #include "thrust.h" @@ -18,5 +19,9 @@ void scan(int n, int *odata, const int *idata) { thrust::exclusive_scan(idata, idata + n, odata); } +void sort(int n, int *idata) { + thrust::sort(idata, idata + n); +} + } } diff --git a/stream_compaction/thrust.h b/stream_compaction/thrust.h index 06707f3..90889f4 100644 --- a/stream_compaction/thrust.h +++ b/stream_compaction/thrust.h @@ -3,5 +3,6 @@ namespace StreamCompaction { namespace Thrust { void scan(int n, int *odata, const int *idata); + void sort(int n, int *idata); } } From 3929696da2256ae90d7490d4cd9d1bfd024a5abe Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Tue, 27 Sep 2016 17:36:45 -0400 Subject: [PATCH 08/10] added readme --- README.md | 91 ++++++++++++++++++++++++++++++++-- src/main.cpp | 2 +- stream_compaction/efficient.cu | 4 +- 3 files changed, 89 insertions(+), 8 deletions(-) diff --git a/README.md b/README.md index 2cb2434..61260d9 100644 --- a/README.md +++ b/README.md @@ -8,11 +8,92 @@ CUDA Stream Compaction ### Stream Compaction -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +Block Size: 256 -![](images/naivevsefficient.png) +Array input: vary from 218 to 226. -![](images/streamcompact.png) +Max size: 230 -![](images/sortcpuvgpu.png) +CUDA Summary + +![](images/GPUCUDAstats.png) + +![](images/cudaCoreStats.png) + +Memory I/O looks to be the bottleneck over the GPU, meaning changing the variables over to shared memory might make a difference. (not checked, just speculating) + + +#### Output +``` +**************** +** SCAN TESTS ** +**************** + [ 38 19 38 37 5 47 15 35 0 12 3 0 42 ... 38 0 ] +==== cpu scan, power-of-two ==== +215 ms total, average : 215 ms. +==== cpu scan, non-power-of-two ==== + passed +==== naive scan, power-of-two ==== +total time to run naive scan, power-of-two: 187.5 in ms, aver: 187.5 + passed +==== naive scan, non-power-of-two ==== + passed +==== work-efficient scan, power-of-two ==== +total time to run efficient scan, power-of-two: 171.601 in ms, aver: 171.601 + passed +==== work-efficient scan, non-power-of-two ==== + passed +==== thrust scan, power-of-two ==== +49 ms total thrust, average : 49 ms. + passed +==== thrust scan, non-power-of-two ==== + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 2 3 2 1 3 1 1 1 2 0 1 0 2 ... 0 0 ] +==== cpu compact without scan, power-of-two ==== +182 ms total cpu w/o scan, average : 182 ms. + [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 1 1 ] + passed +==== cpu compact without scan, non-power-of-two ==== + [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 2 2 ] + passed +==== cpu compact with scan ==== +356 ms total cpu w scan, average : 356 ms. + [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 1 1 ] + passed +==== work-efficient compact, power-of-two ==== +total time to run efficient compact, power-of-two: 186.149 in ms, aver: 186.149 + passed +==== work-efficient compact, non-power-of-two ==== + [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 2 2 ] + passed + +***************************** +** RADIX SORT TESTS ** +***************************** + [ 3 5 0 1 0 2 0 1 6 1 1 2 1 ... 0 0 ] +==== radix sort, power-of-two ==== +total time to run radix, power-of-two: 67.8532 in ms, aver: 67.8532 + [ 0 0 0 0 0 0 0 0 0 0 0 0 0 ... 6 6 ] +==== thrust sort, power-of-two ==== +515 ms total thrust sort, average : 515 ms. +``` + +![Fig1](images/naivevsefficient.png) + +##### A comparison of Stream Compaction over the GPU vs CPU + +The time was measured for 1000 iterations +Exclusive scan was used inside the stream compaction over GPU, which is shown (CPU vs GPU scan) in Fig. 1 + +![Fig2](images/streamcompact.png) + +##### Radix sort +Implemented Radix sort on the GPU, that uses exclusive scan. The output shows a sorted array. +Limitation to this implementation is that the max digit that can be used for sorting is 7. To increase the digit limit, you would have to change the lsb iteration number in sort.cu to whatever bit is the maximum. +There is a comparison between thrust's CPU sort and this GPU sort. + +![Fig3](images/sortcpuvgpu.png) diff --git a/src/main.cpp b/src/main.cpp index ce0c6ea..7038f23 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -19,7 +19,7 @@ #define ITER 1 int main(int argc, char* argv[]) { - const size_t SIZE = 1 << 26; + const size_t SIZE = 1 << 20; const int NPOT = SIZE - 3; //int a[SIZE], b[SIZE], c[SIZE]; diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 42fc04b..98a22a1 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -23,7 +23,7 @@ namespace Efficient { } } - __global__ void kernReduce(int d, int n, int* idata) { + __global__ void kernUpSweep(int d, int n, int* idata) { int index = threadIdx.x + (blockDim.x * blockIdx.x); if (index >= n) return; @@ -56,7 +56,7 @@ void scan(int n, int *odata, const int *idata, float& time) { cudaEventRecord(start); //up-sweep for (int d = 0; d < ilog; ++d) { - kernReduce << > >(d, off_n, dev_in); + kernUpSweep << > >(d, off_n, dev_in); } //set the last value as zero From 931983406e8672de93de8e0b3d9aab53e27e4328 Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Tue, 27 Sep 2016 17:37:14 -0400 Subject: [PATCH 09/10] added graphs --- images/GPUCUDAstats.PNG | Bin 0 -> 28590 bytes images/cudaCoreStats.PNG | Bin 0 -> 44885 bytes 2 files changed, 0 insertions(+), 0 deletions(-) create mode 100644 images/GPUCUDAstats.PNG create mode 100644 images/cudaCoreStats.PNG diff --git a/images/GPUCUDAstats.PNG b/images/GPUCUDAstats.PNG new file mode 100644 index 0000000000000000000000000000000000000000..47ff1b3309016f9f7717e8b132c13e5d665bd1e0 GIT binary patch literal 28590 zcmeFZ2UL?=*Y6uZL{UIDB3o&S1qA`IfRqps5er2{h)7crQA+5903j#}DgqX&l!yhC zk_ZT)1yGa@K`}%KNoXOGkWfMrk`Qv@cJF=m_r34AcieNwxaS-9oM#LNJ{gNAYdveO zIsfZ7=i>fZ2U`XCt?~c>K;hKMeG^& zRlZjCRscXn+PdW{vXc9?!6#kA0070-pZ}!VgWg;N0D`Pf9k;p=<;lc)z4x%mpKU-c zi~`3#uj7p^)*9*$dv9a5Jlx2Y`D*JEZgLy?p2wYY^lDiJG&doJ9pAb`C;qnf<@M%> zdu~?SPfVs=UCQcT1>NzuL(A%_+EDiHt4ALpYBpHk*>+Pzay?u+RraiEYS_OVJ9yP) zp8s%uOPO=k_1==hzWVKgqB7uaNf-aR)*b^W{`G$V;M9M-TS!!B@!j2iEnr_*Xj^*Y zxuG2c6MSDIOKf}khj)7l8s?{Go+9O9ihSIw4;mYc82|O4`2F6-%K?!~Qf?dNSF1Z% z0(U&+Ze=TIo1eRr9^v-1p#dIXU2GP6oxScyehm4?H5dQd9+mjn98L7%Ta-Y3Bsypb z(JJ-V({4QHNJX(m6FMK40fO8Bo{Q}w3%45S##Y}VrZ7P*0N}Z*AISARD;32ztqzb~ zihCSJeg?02NY$(!PPc7EPNsH>_d=gRofU!Fl{ap6I5!`VePFH4z->5ZLtokgrhfkG zo8|-)fw#^l2DN&vW&EZn+kM$ZZQ=k}rzg;w=Yf^w)3FC z`3GWiKFV#!qU*&*ev!N`?!7Q4*FOVy&UNa%29yK)kr;&Bq}b97+`I4&rjL8DJaAJR zcg@iGavbY@ha3LE{L4zSKKX?g8cyOTWr1!hGZ;2OpzhY2>7xl5Cg^>wVi#309%>jI zD5%EP|2RqZ;=$Vi%a1yqLGO=^Ke+kTX0d+pc+0e4TmneY>CS`!jxerGga+RkP!3iv1zlpyeBxHlK3 zVdX{(w-GL_uBXDg1=}pXUo*DTRNmH*zpWo^q}=6b1nM59%b7W=u9HgRZ_kjp&gGe}|hPL-N8!Xxt?wNbw_zy{3d*5!~=Q9Y`Or%=Vb-Ase zs*&+9`Pj(M)yb!O?-bNQ+Vk=2EtIy?k%7g&*b{U7{o(R0$++R^m~)2v3k;?%pugk4 zM@B^a9P9kY>4)Qar_YL*hy&YUiDoo+wRkXBi%^JhmGX2a!q9??O1NB^;A+u*1*{jg zcWCupfcl@a_`JGDBR=V)kxg}Dp{I$|4aw|nJ)Jpg{mg&UhbymxY=j&@RoMRN)%PPv zh~MXKtPd*{h0F9s-`KjmzLXXze@0;}<_r9!o4(RusioX4-1E1HXQM}sjX^k?!y!vJ zFXh4}s5whetq%dgZMVj&W;EzA{2Vw5RS#k|N!2#wpV|lk%gn8hvB%^;Nd&kD#$3J4 zIo6;c|Bn@ya}Z$It@~m8q;#-OlyUomdMn12=kl6NQ=hXQ3R3yJ!pjFXj7JZCJ=N-7 z7(N$;l!}_?o4D)kh&3!B_^<3M;q8s0+{WDs(<%qXUH?Gq_dMf4fd?u=UyymMb|yZ; z%=KG$PcnmdU@7PDmudc6cweDq1bB+ufC3?7J}R{22Vauyk0d_5B|ED7L({N38QUQv zRomn+SVO#i>Ii`6607#@!5pvTUG6_$>X!J1?e1H)H5BH&WvqJVgN=^Y&{mfh6y5}` zXQ}t7aCYn1n?pp=YcAF>l2Z7a_c(ZsB!qq0+j8NKC9#EegaGd(LYB&^qOB>?pRu0$ z*&NAhP>j3YCz9vB9Z|!^a?OTFc^Go$f+v z%fR&6&1nub&;j#rYtk9 z%SLvPc$G>%Q6q1?d`@m%AEx&qzzMlEflh~)UgB>P`4$TSChEtbCJ-A=js~(|FY;}C zpkj-8N<$=RQ?N@B;U-z;FS`Vczbpl~pHn)Zcjcv8NAJ@1DfjJ{4%_W3=j~`;JFhNc zEJbWhgULxV>)WljiyNiq)*H?_BX0s9v5WzPZ-VdJYXjG}sNBN6I52oQn`i?d{~o7j zwrcbPdqj|;6`DV_Et(ZVTPh48n_%|E8WOj}>q$@a&5Z5Hn7;WK)A*XSq9t}~I<|}u zgSrpkXP*QEST$c#;&T6vFVSPNqwE)(JER`QJ8U=l>tjG;;dVgMo4w%oCZgt_`r+3# znj$gi<0s(_iY^G5zdk&wA{~5h0|=nA{8Nnlam8;)SogR4k$iG#)!$Pix!bw+|Kqz; z1gRw@#KIo!zptyMGshjJmJlm6gkp~sCg~pY@MnGj9eS~6;(@5()tudHGs6uVe_j>X z63UDGAEp4re9LmRhT4UgdwXx{1pT=9aF`GA0;jdV(QrD=JGF&f8$e+)=<_vdY~P@W zS*c%NHNNM4vyJ(eFhJ9+5}P0@)G|w5I^oiroIaJ9Kg*5aL_HifdL52TH%jZ~lA2JO zK@CO-t_x3d+pk|wJtlN}uMCtMZN_aaK(mxY{UM)BD*ZfP9Wt4T#h^>uy}zaIeJyB){?&JicwspL@gpg9N7G|xw)X3~E3^i<2T=26q}^VU zj!(}-CH&{ph>#v67LS#?_v?Duzp|NmTuW*RIaMyU2D!`4r<6QX-P@~R{>Mmvul7DY zZMoN5GmH82VdO~fm6>uD@aXYBF7>-y%nhAr#TKtwe}3q4m5H=cp*1L6eLkh^ z|LNB&wJ2Rr!8y$)#f=^rS93jVc>bAMWD}8qpD^-kPyY%(-)3PzgbM4U8<}m(LvEy) z_+LXmBF_x^;;^V5FW&AS)l?5!v~J%>=_r9c{yOrG{d1zy&Y2gbvz{K93K{z2ae*dN zI5M3q^K0;*&fP#HN4}r|mZ@Ozp&*i$fHT4CusPHBIC%wc*bFD72jp?OH#lABO`aQB zx^I^ceQ7awzB=}zr2)b7-8l%%D7-9&`Ns@<#f`2uJ)FzdQAIsb09l?zn-n{C@(e7! zTp)ap(TiD@k;P^1nY@<)w$(AOV`pl?%LK;3| zLxlF+6;poSU2eo{&zC|fORp%}f@WDanddf1sT3Dh`TR6KR|1q9lhp}|o!ZUs)cV*i zobN*y7^()um`3dj2oop3#eDq(6_W?*rYcUuxUisn!RgoR!mqbFmWRC8SYysFo9Vs_GBzB<3vy?gsD}+C|0BRLANb%Z(~EUC`IcOzVzN z4M=zMTCc(I{AHwI&04Fz#s%L6b0PUozivDVuTjbKgykqwykH#P%CY6yF6_{EUO=6D zAA0d~>fTsnx?=~&tM)Wx#p6e|zh*q>9Voo#nMB>0ff{46vCKlSA2lC-8il>~a)f(L zbg?)D7{Vs_gi^tk3_cY{iNWsk2)z}{t)WCwJ%p1yYE>!gN&B&OGk!u)rw*#K)rCWr zEOjn@Gb_Wt582fgQlMDY8WFun3CcQ17|}OOOpT~46n|(EFK_1ebnP|e#W)J6j;u$9 zvEzB#0A4ERgzVoaQ>g{w6H+GLK7he;=A0&;U8@;0bj+hm!08zA?PI z+A^s0u4L9B8maEoZZz=)I1Smj6WtJ)|CaXtCDDXdkmunGK9~Ak5#P85mxAp0<{OkQ zr&Q}Bd8tTofRR>Tv*@Vj%1g$iw$Ln0NEqy-g5b`cu<+cx!x2>h?`g4ZHyt&EwK;0_ z?$y|HO(CS!{CEf=f_#EmPkH3VqXsUIpK!y>B>!}LaZN83TJ(I13pgHiV2tqK(A=p3 zL#N1DXH}|H-CM-0b5>`foZ>YxX=g4nVt7B?ls|4xH7S_^Q{`u0^n9s#iGC|RXIm*K zWpHM)Kk`o9l%7+a5hh(@zN&*v`HFoIuV=a!s`JqutlXyav2i*Gz5$F}^IAHXqohSX zT*)TnnBmBVrA<-(^Kfo1)1*USjwGOwLm?@7*N5|JrzrMAylW&L@*p!kaw|dEBjLE6?DsRWP^=^~?s?X;0>RdI0huOoVoMv(a zX=~m0xZ=`T0f%cYFd7PH)lF3iik7?eiULIC3qf?mT&3C1#X~+}qvqm^>=cLHEGIUs_YC!81Y`1}xhJZc-@!v=>LIwcqUa4GC9a5u>9Cu)efZ5DY-d+E3A!+679(EjIl;fzbYhB+>3&($(-=p6tFdr5BTfxwo;&%>i;9+u zQ2G_-Hrf`7^oQ>k$jl`mmf4NJ)|;LxdPEz9i@S&jqnZvW%{Rsp>klvV+CHOSN6(I= zi}p}<62CEQYdWrRcK|`)YHyq~+Fm29Z!#h+r)C4O@aNi`e^kbwNm+O*gAXkYTec;1 ze#VFkj}%GI>BR|3)MFPFO&JrVY(o9Aw5b=NyH6zUS{Tup6yOjpyPG)oX27#m#-Lh3 z(t*~TIz=J|xN&I4J4$uLyS&gF^vro`@h=h@b*}X%t^BI-=q0=5J5Um`OifU`_QzK z7t8YU>`G(uo644EDf`lUr}_)nH3iC=jUB>@QnWEeIYM1lKroIbI;+hj9M@rA?zF;s?? zx>EfOG7dKLE_iakrA>Z31IYqTeLt8HcHv0eG|BKf8kvP~tvb6W_K?mkU{7$L0q9PAH7Suo{=k$wI7<~q~VDj%Y07m zLYQ;RXILAD_blvB8@eJ5QR%3o^4xW__7zO)QC(E#uKTD439Zlh3aO0A2Z$JFMOK}4 z-zBy+KR*nGDjq&NR-;cS0mOQf3*wL-2yDsAkDX|Y|^`$VRoHk zNEY0NFtj2>-E4+W``;d@iqO2=RCW>e(&G8bl@@iv7}rthIvUzFjH_#}gA%67+E zpPc+Q&L@3XhlnU_-r*bgthejJ8 znbcdHO!$z$asV+}@)B!N%~MHci}>~cpmMw_^-?iVWU18R)i1S_;Q+J24?YQ-j17ZJ z^npJbweUf05{Vw(W$4oNO@7vwOWsQ>YqVksv0*(+-Q`)`vTkhWVH3-&tRVtDTfQYB zu-+$SN@^M2+<|Hy&YM^q_(L)@XU@ir>DF;R&q%tdWpP%M67o5^Zn;>lieTd1{&FNW=l$m_=VX4|U$?Q|RRWm4oxNJbhzvq#)bvJ(AY_Av;HMH-B;(vQ}{wK9grMc$P+ zsatb;M&6sppj{B?W)o12Tlo+2P6`-p576}U<_hfjbtSQ#rqK~@o=O$)64=H9HM-Jn z@*8mOW*^?v2Tn4bV>tV$$};YC03ZG)w%UEZ3d|<>+trzh$IVwBVajWco{N~Nuy*W# zUI3tm2WG3x82Gx9<*EZAoE@Rh-?$n1^VWxr^g*39)d`LZmtjcu%C}vmlX^-eam**v zjs-ht0q_qXTmDjhXy5DH$#D*}>tNALoJu?>C{@3~0>K>gx=xJfk|87TiH!wKi8n!v ztVuXsDO`vg^6wLqL?&IAoSGM5d~#NE0Rcw_#fs7!DY4C^h;J5rbsl0}s5qi^ ztH?yN1RW?@0dACfkw2}Q)+~PD#I1FQu59JEp$I*Ul%cBkry%B!v?y5q0e;fk1|P&Z zUSk)(Ce}NoZ<|1iF4Xc3dR)~kcy8hARVua*r!WXOX^oh1aibQrE8CjRswjL605xb;`^U!nZ+OX3Iy? z@x79pN-a9FcBSI7n#3@Px?n&_2IOv5{w1H<=3BgibEMP#G54lqWX|i|Mie;BXva z?h__SLgPILEW6V{)Z&|s)c5gSwt+}1}3V4JQDI~K9P1~!dJ`HBXf64 zuXQI+VHL5=2)rjlF1a35KhCxpDvgS1>T!)(AvjW3)&XdlXVWD}|gzTOv2GkCl_{ROyk+sb%@KUEaH zK8L}2oFEwS>cnS}N-(suv|)+Ds1c3Kgk{&Bmcv8%%|T6emh()~LM1FvW_CGBGn$4Z zyG*>fMQTwIce z;lFL1Ste4AoS;KnVBeWFVz$4D3RB@f`4S0l7w7yE7Pltu1zp28p1cY9{uhxMU;qoW z;P9v92UFFQDYxf6tLfh)dp7ac6Lq3#K?T9b9oN&bi}DPQc3z6^^tffV9T8Q%oW|gG zlultKu91n)$Nbz&mOTzyUal5YNZyuLQzWJ_536}xY%3UY)2g#EsGHv_iaz@B%)6I||KIc;u*@eGSIVTFzd`5bRD~oOenrggCI! z!SSCw1HDVXfc0|8U$L*7Nb|}0$yx@P12N0M*5)}8_#@{4VyT+VnYUU*yhma=l9w8p zVki3X`4;IvZV)9b0wn&~wR|nXhktg>laTsPhSkd${uhRI=qJOn1sJdi!a}Vh{~%aX z5zPNUuo_Vs|3a|XzX%pWP}SqcON#sPKgnTC|6g&K*bTLyMX46szsZ~Y_n+Kt(=YCp znvxBQNsdwa1up?7tW271Vkk1PM-}UV#{)7!pvzrN;%;i^%< z#{bjchT+R~l9aAbK!EgSZR#;U{1G#cr$cDb*1s#i z>h3@G3LM#GW1H(6Iwm(M(6O(ZdD^QbpP`g^<)LpR{f22M6Lc~mgdXU1t<~9Xen5c`5dGebXk|(i)3GtsO;_ierjy zhjvvZ^~9Wc7n@Xeq{kf-8GsJ3xN>WHAba1lXT2ww@1QU9-d%nhP##%6CTPuO{Ar~5 zzgqzyx5=HSasna%mIeb3G_GBmFNmw-bFt#RA7TW%?Wodvo#W9qk)DSPwBAi~n|doJ zWeCa9m>~9(7qOHnrOsncmU~v475HMhp6qCw{4`@un>-7!G~Z>bRd;?NHpZnW)%Xf; z=J0_)>mYnvZTWzfN6Pg5pzc#z-PW`x!pRXR`iVt*pQo1R+Ui+zFd1lRpj9Qc1P6D$ z2cm}6Ftp4cM~Qcc_9IYML-*xnja~whhdsZ2c3$t*!`BM-)CUP}KH6&R`Q1kc3m2{w zp%hb;lyzvv>gnamI`?w7^N7$0?7Xt^Y4U|3h!gCVc~Mc^b#J%&=b;7KXUQI7`^(rv zLdwVl2~Z!#|9a`R#xYrH^0ltKM(@v84%FPx_7D2Fwj~$!^k9G8pXCJovM=4F3=SqJ zwctw@jMUs4^hyGQ^S#0@p;>)hi*-XWZn$a3!vOFb=P{d;%K|@33Lk%MdZ^>EFboj}G1+Jw06%s228(hIp>w+wlCBfFw>UJHhz8#TH?R4kAx*_{ z5J_cgR`pkhR-ccV9xUULHuyQok|6-dS}(+&FpXbU_FQIu^(ejBUMg~<9qmfU<8zF& zLrR6+`$(pgt1>qYAG!p!_h#-YfjdE@e>4Kk2QAWh#};24+}#lgQueju%ZCpXU$ua2 zC>Z0SZ-7;D+2P2pkPK1c_bs0HUb)=T$`>bedS&;!$&z)BNY%=J%KD@hFVt&c2RfFx zy}n?*5SJZh_W@7gw78;|25-F|reAZHo2oWT^ouJ0LdZ$)Z0!ZzQ7LUW)3}Z7CRGrG z9)~0xK*!+O^q}g`9kq%W(ov$D6zVqZ^@@v5eOdihb8kyHfp7sy$!>no@YemD!W7A9 z6R~>JM8{2C@ZzADB|V2#b9{8a^OUJw+V*#!?;pE6DTumUsw^G+r_VMGSv@*2;knJ@ zPA5umDZ}68_EnV)L`=T)oFHWM77a%Tl0&R*ftsFJh(FGu!>QYYW?^FuN9cQVo zRCGd>0%W6tWy4RHP$5xy*lt!KDMt8iCZ<~RP?p=%^hF{R4OAISldCxz(xory{p}VP zoo7d3L84hpfY;d>eEG+j`K!rK&4%jwHt#W&1gbx-eu%bW3#sLj>?ou^Ajb;;0pEkl zx3C|&eKL@n>W-|J`>F@MU-B_2NX5ZC7xv(OWgX&rIv^r<-^^|voo>y8jXY!FobyiQNHpB#Vx~pLWyOr zceMX!I0FAWxcqI*(YH5xQdM;XqoOdr`@-FtFC z>-C~EL!ec7p;oVjL@@Aq=K;yT+TAT!sn|-6++zx-=mj+gBwd0XX;DNK^qr+rBlm2% zmGM0@PKdhAkY$|?Vz=*Wr8!H4K&llvULOv{Xxw{Nc=x2xxI5_q?T6u~6`&oKF1+Cc;jPx%X1&GH8(^vD6uHk^ z=Rt2{ci_1)+s%5d!oq>eDm1De~b(<2_Tgs*Oo5FG>#( zUEmqf&Acw6!KXq`vJHTRQbZ@V+grt)5aV(#r38YMp(@dDMrNDkJJ}QK9vsjZ&J7#E zw&r%}OJ=wmU|}7Ua24ebhN?;&Nr&op%Z&CjYyM0@KPNrh513Vm+o=RliLV+rKXpD= zOQPw>ZLeIeBK9@zI^BJb@A-Kp@AZDi^Lc;9*M7*F7NE7#+PcI_%v}a8XEFny-mjometxTbfz$r5 zJZpY?U;EGHq?x`h^#G3j5%DNU^i@-s?=DO$T=;P0;k?CrI&&dQkCo{w)6OL?oiiQj zamA#%tzf`FS97$gdFu3+piWA2XKaC1-MIoZ9%hF!o44oBS0J9uF5AKacx}AST(dm( zm>So)>x@5bCD_M^T&cyXJ}4IaCP@ZYrjTdxA7(ros{$Mwj*`PZvl178DS`Ve!Wx_+ zj~XMd_qok^*2VJNsccx^h*2#(9Cf(ntYw5Rr}4B#G02eo{!^Zb%ND}+q2{n^^{Tmj zIU!x?>|7mt?;E;21>Rd;g@QL~$1h8X3|*31YE`=yxKSRuMc3DO^VB?!jsdGkr zqE9#V(^&v2)zH!{ZZ39ZrO;WYvv@XoyP#5owN*FGN1e? zC*E)*Glnpf)a```DV||>;bGnjgfp4llznSgje^|wBA&vv;O~FN330BsPX!R1G-vD% zZDr4Jwq!*YVi=9H_*)(euP5kZ5bwHBihgLcam~bo{f9T*gX-d#m|5LrkL3bA>PFE@ ze1IiTDTX&YvP~`EM?Po%+ohu;{aR>L3EM!qfOc!)_C#j)^ex;~OC(nQwd!}HH;mkt z`kMjNI(xW$TJO3pJP_C>PP%ogFXGRt0t59<`ML8dzbowrGzZo_Oc?@Pst$B{q4I%a zoO7|!YLKq-!!^g!Bqm6bl!(lUlz8`{h|Or?W=L~w+*{%xDx$qBn;BF-3yVo6FeCe_ z344hBOwv@@A+3)fz-4zuz2oP{xCr3v%CabolCf4KfQB+Fah-~8j^)33Ei1`*$Xce) z&$w$SwkWZl7-~2t5|zp?p}c0XA+KXSWygvNE?HS3ok#7x&zHVaC43CDRk1VVf;|YS z1nFQ&LQdNxjtNhZTJrD}o7Bxs>2-zS>9bvabIo5wTOojzaeFN#)B%Vtrf??F4V)zv zwLgG$7BbT?>)GIn$dpBGl%LBLyrdh{Sh1~OKfXE&3d34Ks~5#XrSAmayGnwT>!AmY zeX1Q5yq&BX)Er7c3Bgu_eEJ9#U%9PHH%IqaYos;c+FUZYy0OD<_ceCam9%kGuj z##8XNoNE0u7dDi;z)S)kpI_I4r!&1d4?V6ZM7ad4 z->6W=D(6 zTqEK$$645(PRREU7(R_OBPE2ErpIz>g>~zF{$7;9;gWdyAipKrK;p0%qr!`>?qwfc zrmQf6jh=7N&WRnL>mUFm=PeO1qEO4WT=MP&LvjV5BPxw6#ltnE=R}78)Tcm}e560w zvpUG9#WQnEt$Xaw$(7A}3`Y=!`&|AIv@veaEq3KP0$t=NlFfc#e;K*--4JEMoRURU zD#C84|1G|uJwZ}N`L!VwMOGe(rNl-Jg;q(An1XpY3d+Zlr@QGls2fTKOyd}RE*OYA zsi})$a)?koq{9B5{Ox2svCgx0=KFUfy@WBK`+>bdI#}%lLmcQjXOMFQ=p|)I_%1I| zpi5OQXlB9qx+<~4HQ7I@O;c0qP5|K+TFnA_Wg=}-8vwPwqEp75*#(d+RF*z~Q1ZwN z6I2sj!DNUt;cMl-)ARh+AkP3&os)jkcE%X*B^OuAK9H|bHK%e#sEv|hZH4!jU%l>A9o*#JBFUsP+qj{I}8}hVOev zn#0OtKdu@r3?*KdXxY`%oOdRh$a=W~CxvByF<^f>VC1l{hvP}fZ*wu#*$@vElu2@> zkqGukVb5PaLQxR}kO(J_;iVlziQnnXHb3&w1t>2`#iD!CqJxLwd+>*vFuS|d4@PgJ zr_@z0l|J{BBbe<6x$K@XE%Ds;j!@b`S}~-*Y>+f#if&1 zfA-o!7JzQwEMGNh^E^L8i?bOjyR^}VPe~a~yBsYAZGvK&DK*Wyja!d`!66Ox+aoKO z_uY!8-8Ei;U0#@^^Wf-b+*VXZlJ!NE+RZWXF?t#KkE-8}y#EqJ!K#qWvT-uPxh9F8 zRTI2y8xK-GKeYPwZ^mwNhaE@e@3e2OmBU*eepvnnIh}S8D%q$0Ds-zzu$CzENA9j4 zXz2HfM9I(6*+eU5WF!n5QT`xlzDMTh~ zO{!8IvMnFq>~7#M?+;ActPSLO@NI?5DI%~b&$h_DTwvLiPQ-qY`*F6ds%ksH1kH@y zT2+VvuWSx7Pmh~9v>U7Npvj>Qdrvi-5qNfCMHGSA9(g@TV$xredT+itD@OzM9y_%DIHM-nwe?^6v(K&ED?`CP-PD!1q9Q~H}j(7F4@>l(~ zpMOwzG06M;mBp_D`I*gI+1Ljn(=%5JUL7I)`O-+g#Y1;U9jWR0cI?$A>k4w);tqaK z{cUK>>19F3C{#|6b!NJ@K*#>u;2f*L~ZY zwuu%~Lb+m;rpJwW^6_8>z44cJkLOAbplsOwe;RNP+y=y+*xlZCt>=iT_XFDVdE148 zn_FgU9&2Xj>=Pa_HZ~ZD>A8jKXggx556qX|uau~%%_Q#Ay0j!!w{nW@$?cp{%HfYm zIIUIMPfmL85NiUHGU-Rb(A@taDvBBOt)qOOU5Rikv8VxF55g z8eObT?XdEnlongG{5-(JPjEOgyX^aw2kS;Aa+1n7xKT$3U)%7ne7@hyQ4xWB?L|Sl zhNg$icO}~)0zjlDw{D~CFkLu!@tasfZ}Odu(}9?dM$ES^Xvc%ryrzAmXYdD7Tu8d| z=6xCEEGxXr08D|Pwdnrm*^=z!FDS{e49@pg;mTenDyr$kWlQL7iO}$!A?rPO~uFkoo6mC)F6k4fJT}7}3Cj_AFWXZ-m z5iT_brpFu_1veuulfw@d=`_KgG`?N)O5d>qBZif)#cHq`&KCuaDN{0^w%P9|U+@e@ z=EX#9d0*-uSC&~{YG!PK^K2&8jAham!_XI=o5y`T(K#`eP#H14>YVF*;n2;x7ab@} z*Vjuwa`E|Pav$P6RE}u{ePkyJ6626nQ!7c`>lA%(g!v*6o%15`QS%YQ7;L!RW_8eB zWv)dMOVA89E_ipeEz_ZEavlO8U1+j)u`|9yPSns-d|F|nJOoqr0kCApVZ<|>TuJrZ z`f${rwdCO+e#p!jv@+VXIw5MhZwy@gXC>Ib;X!ZW3~WEJ91QjK5sb`W%V~E8JtZ{{ zdc$j~a|>Hzoh)-i{g^->cxckab%V)aBSw=x6IgxoB$NHE>mPd#!5-%Ye|Q++c~xc( zUI|-}hMm>wjq)jWpS&}XaQV?hPX4p>1FbHF+**vhd-;kc)CFG3(>kZ~=CTC4`Rt#E zLcQgaz9>l5CaIloxYLC{Jv*=cDQ=u7-skVHV@4xEQFdnV^Kf;j%Pn36y0N8x`>V*- z@!N_K?SVI9PhKY=&#msTUelr`1hv;H_M);5<{fu0UycMv>s`1P7kgMmDQ+{C`-{eM z7V2c?)Kvn^rD`vYF;Rj(vKHuYE=l3ck|5zPG65j%e*%cbEkd`>6;ptvN+#rbbJ&&)|T|%3T}Yy@qAxO$En0K{J}2Gs!VDWLHb8azeK)rI=$Y! zi}BhLFHx_nbO6vag?k@TY-Pz>A=1Ga{>Yi2gcrA2S7oD`=x%`tmA0~DPoIXq=gaLy z_@|R^YC6?%PDaE$RW#hsIzlj|t%xWV(GOU3!mKVJaqki`RpunL-T)<;H zyyN~0(sNyzd!jRHG5`HA&HIEeXo<*aQ8=`6(k1X=8hE%ffwe5xcuwwXM)@h-TUBZ< zF){I3osqR3B)MV_o|$Cw1h(;X`Pp%;)#{`$sb{~RzBgGu<8TCk3+Am6cT9FhOE9bI z$htOB#sv0jZ}x@n?vnE=HnN13%{%RBhcK^)_sFn>eL#y1C0T?mCjhA1l&hs1f;f?> z^3AH7-0mNG-WsTU01WmtRJcQjz^$7l=T^itO)w2GVotHCK3u6Xs#pkbTBj=&)kaRy z>5yhOdz{hHah3WrHBo-;y;SX+i~&-4C}jDu`>Iihx$YZ9Sq661XbljbDKzV60c3(} zfVbj)KXLCBmm4WII#DVf^||Es{6d8k4DO#EMt!nfZuGa2Dk7~ZjN-9sRJT#Fd~1*{ z;f*sHvfmRPwCDEx%+h7d(Q`53Y35e%gqGqRPpUI0R$bVp($dcfrS18G%$0(wqLta3 zdkiOL`rMv>W2iQIN=0p(G?)(Gr_`#`nD93URvdSiLj86fT^@VpNp$vzBcR~isu7Lh z2$`ia4d}Y`tFm>hmv3nt!gG2g9|ovh4>}$ShBa%O8*WXTlE_3WZa^LB{w0owP$sKy ze&)FA0Tv3t13%%NHt~KICMzL$lyo#1Ro-J9u1ok(J&>rao@uVvi#^-aB#lMJP|N2c zU9N+{CuxUPwff|}>0$PZYGQc=OwBYz!4d7m&QH)=Lnt#>-VKT2^cVrRql8pRA|&34^%12Af; zFhX(~E0J7qTx)%G>4W@`I>f4hJ5%)*dbyFqoEi%wRW8PA`}!o-?{W7oNaUVGE)H81 zN2kqwIPiUbQW#AOqkq!=WFRx=np5c=lBy@+p4Tl0KQ3tA-1=1hIosd2>7~2Kqik{T zRyEUWi)9SWwEbBZDT8Qs1JlktjAbmvq%?3!iLV!-=nE%3JoUU}%Wsd{*tkNr9p$D{ z%bjuh$9}X4OeNu~aK$7;DAn|nk=_FKByVW`xJM>fB|3K{EhS?+S8@z6y16A!b=W-b zuOnW=A8tt0BLMV7>e;D3AYP2xQ%ROlOV8ydCMSE=c?Rby&))g@uM`?H5um>9*neQ6 zSM<_1Z~TW|{B`Q^IcY<@%<+@|PA~q_RnM{i6&kGnQ!kRzB!@qDT#f#}khcwIT(G1?k^aJTgXfw-7>wp;O4s-8Rw{sYLp!rl4|tFq$N%r#M}%!*Gw@$Ub=zGcMU|BhLB*H`@#j~ z0rUhnO?pGhyV$JQ|2#ot~GEK9)Tg)xd$*&g? z#o62rs@OYFzYz}OWh@kY_IYj^Hr24af2pr(vUVnBX7SR!Mi1k%s_n1j52lsh=^#>_ z+I=X!#o%MtUe{0E-c07ZL)AC$5K zT?hNj@P=%1!->~DOviIdArMif3Q{(_4p5`qMFL;iRk^+=*K(8^nbT&GrjuSs2-pUkui&-NJ9 z6TD8H-jA)fIg@mGL=!3lG5N%8^IMU69$ofT>k~xd)V)P5UdK?7v-k%4I)6?f!BvGc zIokG8dR9rI1m!B{gB5+seSFz}?+= zOxWt|43##Ui>_;Zo$JUSUsT>rtZWEzzI{H(Ik@SVYhX%9%&v!Mju+3r&h9&|b=uL1 z64q)J$m6|gSecz2t3t1|T6>G}!eL+^lbJ)0T(7|Kz}L~)TK-GzYw*58_ui2F550e} zCYE|MpvgbKirM*GI=DLHbVZljbQ!9KzQ6vr)2bh8VMtDKJpm*CIay^=cau^e=0#Vs zKfknBLi!Zz1K}?f=B=4K`o?ZpzsjyTDoH2DMOAp5F3y-BnPwXW8e$#=Tc6O0iU8X~ z*I$WvM_-k#`{un{@&)UQ{+zrXR-3bu6A9kv=*2Z4;m+=HU)Q3V?sj8S;d~xkDXES` zc?uCyT774iTdgO2XYqQo6@kSV-;xHS@7|w%{mDrY1-5QQvya6O!|VG-|Hi~nMN@#% zs(~ERozWD~lVre%hiN%#V)eKKEKRCbpCA>n2S6ZT0_*h}Xlq(X-Y(L?GXv+6PU0x= z)O{-*DGQ}76;?+&<`)|zM6gBGytIi&u03xlu9q97c8&QReM><$xu+g1z3Wk`mn3qo z>x%V(a2n}aH)cS06y*8`L+!!>B>XV7iifNXFAd19nw8Y5ctnWtgS$wv+?*Axzr@71 z*0=Ww$pvp&EUT5KIArc3>m0ML#7$(IYV~)b0HQ9C4|X)!RM0BfN*Nz=5$hkDjrC$h z6WQxqBHnopz4)Z?8EYB7rX|9>FH-N1VMyFjh-jmoVMpZ3nknm2vL|0AD+AbqthI4) zo%tos=ttWoQ7F#GHjCrsiRHTct4jlu=Pp!lyhLISUuefR?PWmRH9mr&>iCTQsY^a{ z47n=Sl6(SO>neb}L9iMW0}n_I9%72sd9{+MNHTh3Su5^;G@bV0S@rY^=(nk8lr$3m zk7d=S7EdJDyvs}Ki9GBPmk*H%7LSCwJ#BP2ha!ebdO1E5>l-v1E8L`(oomeG*n0)K z3F=pkelKoSdR&m1qLYDR z?`)9CIf(vaSGRJ~hL0s1z7}{d^5Xi+65rX7x*}C`NGYXTn@UE%)eJRyB&vn0j+F{w ztJ>5RueOc_NfwT9#VrAewJhpWK)x(rY#sozBn#)tdU^pE<^1qNjAp(HD+T9F!}-1Y zi&p;hYw^5opeGx1|NoyjCESu!5G666k2DEiUpCb)G69{>n|&LX^w|RSefE?*?m`J$ z`WKHRGLFj3SuG^{Dn0`P&hGlJo=*R-kkUV~AzngA=8d+f9ZYrf-4!q5a#| znRNBQwnzzz`!u)*r}k7C9XFiLZI5kxFR81u0cAq;i9-$Zt6VR4soM8;6&|f~Ykq*j zgqL3MztEkFLjUk{H+Si{l3a9cFtXpOn#l$F{w!y-dN&L4YHQ0p?YRpb)ZNXEl;E=f zlBE#!qh7qiX?ter7S=djfVWIbWbhX46oS{+m-o+O)u?x5ISUTnYX1akt45!a?vf28 z{y9lfQ}Za!eJTsA9N+pwviDkiG6I4}Y>Wm$nL;YsMZr*7>qpe};>l}jex=Wa9U0@k zqg4kYD)qZPwE~ZYI`cl2?n=WVGs@3cYO_0Agqj27lmqYvAo%Xn-HRM)d*`N;Mc>p^ zjyYt{os)ZFyc6{#KWw@7A+cUIW1)zlW|r2yYT%5H&%)BtrDxUf#Cn%)b3jX>V!@j5 z*!BmViSrlpa36Qgx%NnO6cZ82?4;ACE3Hx4H4)Cli6W#uH6KkO9Xvd06ovFAN3JNX zqjQd98>dYV{{x6wjYus?<}PJublBsjI(Ltczr2= z+T;gyS4I{&H}}NMZp}D=0oefK>$XJPTl^}bbqclDw)htB&L6!bJ2bJPJ2$j9C;SQ1 z7h~`^L?s*4;GFr8c-g{ycI1ri2UyVh>vldJZ(Q*8A1b2J(i2!U+#`-d%NrM7BKW+P zp%cdITp~5GXhGd%vS%9NZhp;1KlX57Z`q>>EWiRCBUIr>m4vIro2#}nN0RycDXnmJW>Pif-Y5@ur=%etY|bfKw&Fm$ZfTAAs)=VawQ9kVjXOr#{qB@; zI__P&u1M-frX;fAI;&^JdRwlp09)q2fgw1D{29p~L77{*=k>wM2-RG=wZ>uk)~$z~ z3R>8VYTXX^=a{S|t6di5t5Y9nTcGcJw5Q7A=MKJX2*I}!w6if%B?itD9bizBAh7+W znAb{lu;$oj2}>KD&7*}Wu)ytKeZ(C4kH#Qx?_1743`$s0SjVx2u~nn+6e}%5kT*V& z%=*-jwxy&XI`JEPtwnJSS)xIH?Te2>A8rFcbM(#{++2QKmM|N##n1(&xA2^pxbu?6 z3GlIL?cYs$>;|P_Di2ikI-K`-mOtNr1tvL+)>xaV7SDK295gA3ER~G=$(Oq)b(iC+ z@yvjF%*0-;|$oh=tc2t%$O^nyRS9fp%!;WLtR3wc+ut)!SUFy{!Y( z&PVEER!J5!#6525(kdq1eh{VLkbKm5hkAU&mh&R-S&Gj!5Qn6bOBJ}IZ3+xo?En1%VFas zMcGk-pvE5R!&8^4jp>o(Ti^$tY?yv6q&8v{pj3>{qlsw=R1y}Ya#!YF3lg%~eCurX zmZys9dV1*6@1qUu&3jhl=^tM+e9SHucvTH>8q~7Xno8p9L-TuVl%IyiRxH%M*9^{j z6L}mMY z3aEd5=-&fqy&>6dEUiwUn@QYH;EkE;jM;8xkk^ZBUkQ#zqb@?NL(obZ(!nH$%fc)_ z1fc%f()O=mp-geigY~d6=!b3XPBYWxw?uVrh9Pp$F)$Ndf z*jHiF$!HFg-{>dy%UJJkv~;=sr3+l^iogTnJ&yV@cm=ZsT-QYe2HB;)-yV%zG~caa zG0fau27C8L+SYAKu7Q=QY`g}tem`x^wOllEfK(f;~6oqIH!>AuHZ(4n2u?S>IHQRt*WW}ErPgKwH;MD)2gb* zWm>mH;!*?=luhccB|<~uGK3IOK}bZ-qwUPud#^cXpFj6r>-d9Zt(TSeeV^xjf6wRl z{eGXh^~{jL7*525@ppWVayJEw&uHWJ>NHY0;9m78)`}ENCv*^98UhYoN2G7NE%cAu z)zSG+Z0nyCY@a7wxq^5=>{p-R*MIN|uAX+|0LmoLNoZu*bn(uQDgFl4k(GaqvkIt#)sBE*3EJh6?XT(jxiKip zFwDvvtz)1%cD zG7szIxNoJar*V#T)3Syc?zJ`5LdN5~b)wMe+pL<`5t-8urbUf9gp7sB?t{)I$Y|~> zk=E~%Or*|TYu?SCKNmH}EI>I{68f=b_>H?2Q1_|WL8563v*ZZm=voikBgvuOI}5kB zeazkCB7*;#qs6Yf)KB*^_m}t4{H0T{HlC$-EQ4_dQum4cx!Gi5lVgU;N2yW2 zC-T~EpWEj*lZ38x{w5h}3Bj8UH=k*U0=h;&vv(4M%*c*sN1ZF3t^K7N(kSt>5sKcO z@l;8&2U}UU3Q}7r5+MB^$(%Y)M?QGrC^Jl1@iAGvYSmEdk1eW*D^+uyja13u=aE)B z9@mlGI+M|vLQY%YvX1{=^KnwE^3?%lGk*RWmX1>u#^~G80>2(wuHd=V%8RZw@*d}TIX$W4h6o9Z~2^Wg$$vN-e(txbI+qGV`qe? ztLD;`Rc#=1QrI{frvxl$QlW1<=0Xn{E=OIMWYud2bP5frZAV_DzVpgO9WU2xoa$O) zE?F+=-yNz=f%>_)Y#xalH-Kb~)4{v-4|(YLJb1VE9BVHh^-2@-0{jpt_C&kIDrKV7 z5g4bqsetODe6w3oMem>H!}7PGk%f@KB-48zCM%8+mWGxWSFKgA-1kYAln6j2bM+I* zH|bRI`7&in{=!w2-_khhSr5Gf9~zw?L@n~;im?Z`{IWi-qA$RC(d!&wkxgrjBh^Yn zVOK_k*N#B!<5teXA@HpPq=9-%1#yxD5W*uZ+eY6Cm!IJzbSe!b4~)T#1;-qIs&g0O zE-su(EO9QZG}kxwzAQpkTb}SkYd1iQUJrdZSkn2!Wzfwu%~;o``4m(px3@%M_5=fE zkEzs1y0Rc^k2Jxxc6z*mZG{mlBWk1Yd~Q@8e~)?HZ-T4C#WteUImHmPEbnH(AJeSV z5KJ=xHVO$0RvqAIdT;T-L_uH4xqgH54k9w@)VMw#jmp7{-r!90bpP^OFEsnt<# zPs*F~204M^^LeG|T#Q(ifDTPly@{ggOzy9Xk+FTsD2|AV6oMJKt9SReT)75w^+rJ; z-Tl%iVdpt)HgRD2)z`6_-?9qj$H@K53x3-GZ4*SJ*xphKu|%;#+587=?IbI)Dy(c1 zcf?Mrx|#moggAJ7^ylEE8-^)wJql*2#dj(x$7`m}_pUmaIPQqOC&-OxkRd_gdV0;ql=g6 zwp;M3)nkv#KFBm=fOr70mFq`F@cq!M5V(w~L>xj_tO~|)IRk;;sZOWh%r`@0k*xZN z;aUbFU&G2Qjyqp@dOCO;Usaja?ZVUwsnZ3$`7gwOSZ@td`}IgD@!<~}Bls;ww3|Ko z#@{#_cc&Jd!dDHiX9_7m?t<`R8{}o)dI?fwKCT6wtVKo@XIZ)H)KG1i;qGr!nA91i zSd?kKLB8J8EZsNDM{q&q`!6|zuH)8no7( z4Mrtcs(q1?K`pk&dSXsfij=&!1*wDE7YdsF$7tIfB-sA2=QX}V-%@lmEn0Y(R zl~aY%^RotVTwjsx>KFq0qHZz=Ie4#Uz!KVpAZ;hwps^pO!TD2dgoxznk3BabzZaA`~gTe9Q$Z3KPCbz4QB=ex=`@SmshZ>JHH8bs)TAehdx}t6D$!bAI~Fjea|*E?5(8KgGPl z!`X@X)MX*EjpZ9I+N=@p)%=_eln!!=<6>3%te#USZD-t zuX#R7XP`Ry^FYueyxh_tIG(QBO6f5+BMuXkEj&Lf@-U{Ze~~fxb*voinM43FD2u|t z=~Qh0uG!`&C0q0FoFso(GqyDBIC$c(*x0|*0l+`9VKoU7%%$D_7p+n-+@xAHGxdLh zi=9sGo3u#e8dEqW&&cuu$v>1`|dCy*U z0ysj~(HMWMGuat?(U%yd+>e<24XZT=q*M9D?&D+8Mf_mZtS&Fp#H0h?N6dZTrFBzX zalg9cMJjHK0QxeEj5<}|j2`qqrwmxxPdb$c?dbr9rgGv}lWvg{n-XD1 zhaLSblwD7pRKK@?EE_wX&p1ywx0JMDnUtIi3x~HP%LiR$8`YCoO9Ci0XO`=7zGC3JQCD{ABw_Dzf4O^Gxa2Q z>%lS@eh`?09O-=MTlwC~Irv~$W0hs+O-@;;^;I{T-NaVbHKP5I{H(yj(;>BA-Ir_Y z-?DQ>Wk_vjrL$xyW#83Q`O!-0>nXZozWpsX_CzS0=}y7zI@H%Hn7 zPmI0rG@J0!bIyjP%^I$)h}|VWbPwoxQ^8Hy@;siMT;}R2h zfnDVDl!-m2=nH}A3{qJBXExfgO8q$%E~~Tn`;~9zlPr@jv@gE;u9ta}<4d}p5Tc@} z8*1g6|DjT?t1JB!ImBsj$&z#+?357wF3RR9H9b8XrrioY@H?5WY%cXy`q9MSp>$WoCv6*@BLTe^SsvWTn^}Nhk9LVOQ3bP z^&cLOZZI}H{J5)a?h}B7K7(Evw3saAJ83=BDs+oz+>+fXNS65E_C(DzWKn^El z<9Bz{(|q?{fxt!y!_=x;x0-Fma(?KL#*6aODFQ4O-C(Qg8gn-sC{J|3wG?EYSgb!S z*&hs~P{)sn`xcYb!u6t>+161_wV8$0EssW*w0@-PE~R5A48PczfSEgW$WK1$u6uBt z?%O}cRgcWUZ4HK-D#d^(hz1PR)WI?Hd+(MX;$E6fl%7*$OZH4&&DD^LvL5tT{yJkb zqQDo<94#$mR_1pj`PpW-f{!Z3V$)bJ#x~h%o;anu%N{J=F4&JF@j3BpZCCfUxH$LD zJm|FubJ-yC?}b5hTHouXoB9)zo|dB<&<>QOJ3vVq&TIWtl6s=$$Da>|7q8oGyJ4C7 z$k8|IH~U+}YZdTZ-+HOt6QiL!F>@XzL_koe8Q6<3Z3~SNbqDqRm|+# zVq{j~ykg%;u7?9Z#(wkU%n|1%hOYR|JVTh9pO5m~7Ke|w* zaE@u}uX1P?E4%XR0lV8ClzgM}cD{GxUF5$OVEfxok=z`M@s<9|Zx>(PDvq!krq*(X zC4q&=n=d6CK5S-cc zb=u<|FW)n6?Z4h`&_epp=MM5N21u@C6Hb3;cMG{A*4|zNsPnq#PCvf7`EUpXJV^jABS@DkYK|dq{w<@`pPlM>@rVcW#hOy`pue^1?=1q^h^- zkt=p5fFb(rjK9>cXeVNFp{F&=0OD^{zh)uE;loWbT6W7e%!qnMHqn$N0(dGD)T6{~ z{^KY+I#Z(f$V-`D#`p)H!LyB>KvMT$;me8d8`MrUpNw0`xnaN_mSDNx z&q@ciJ3(tLYK$`WWYMbLJ&16)loeI)G6Zo>pr0AmE&{fBPMBX^kTY%$?$-8#)I{Ph z88-`{_H`EJr0D7cS0!;1r}l=ILiZ8LXk(1+Td%g~%deMNP<~9;bH4x7u2^DRNxbaC zj1TRKArShB&GJdKveLW04vFKA+!SepA$(@qQkR}oPt|s4~{tFMt=P>KPEgSTJDChWlnGd5=Cp|ZNSD*g#F$FDbuSYpl)@`WzFHQvJu;ph~0 z=%M}Zi;JJ~2ALmy)80zt7)S*5Y^cC!okkH4BKdd)76K5%mmLG^K^wZcV?YsR(3WJs ziex+?W0*@Lwd4?{L)rW@%0iDAAu`2zJyS9-<@^mU%&E==R?)La3{4K)5Yb7_>H9gk z2W>D`Mjskb*Nf3CE2*|q+MtkqGn-l4TVB0kkIt-(1cqNsb=Sav*=iq}!j)}ZpM@*C zvo;*l!Lz{d+xdbxLJC@)t8QZkU%1Yx*G|y-lI6ZUw*k2JY0~!*I|Lf-HCsc0<5Ero zVsdJR88mq5yY42zYQj5BMrAE_gdy-L3pPlY2tEtRtORP7<<@0HMq-cPqt%i24OW`l zHS&d(`WE1B|065SGxLZ3>RFuDsks=+jM8GkFcK+7b#(&%DgP(PKQM-M>V?&ysf(K^WicboVH!^>8r0Av?HX z$ZN3u(fLj<(Sv>c^;l6AJd_PIpE$S!I(0huHLo&Lp~I8xZCYV$67o2wBO%5A0SF`H+x_hJE0xz>#%Y9r>-~RVp zWY)nLg|m?L9?ZnTL(;PFa@lI1yU=#T(7yd|wikADKB3Mc39gIAmIfRUNBV(4o&ZhR zqB%SD<_p9#@R0V~nCx0}2z(1Ic!s%n@>oS%)mL`lLU^P;I(a#vm~3m-`-yMPo}1_Z zNc~qSEF6_=HPKSXIXf$~&pW>Pi(MP0WD3cg-c&NY8ME<{h2Ym|607zaV?N!(p5vm z=dU-fY`~&>H2r}R#5uXH7|T=}{)S?-H?2KztEH;dlyE~V+nc75mM}?tS!Xu17sQ~1 z;k&u^;V82DXap!Bu@9CihI~&gP#Q6et%;P@Tuu7o#~^I7CA5bny6Il&gd7OYXk)K( ztB;iOT+?6TIr25Wc4OGzs4I7{I8n~Ou$LN+^vv=18G8GS7#^x#+ zQ@iR^F3)P<#E0Awq`V+hbUp9T^1AwEMTsvya}!x(i2Z&nRMwuM3r2l4Pzh$+fw-J~ zb5D4Ru?lZqD>&LGPjlN#24$)$;Krgbur%EfDYp;8V=0mr7J?5zQks6b^P?y*<%O_+ z6u(v}8)PK(kGtM|u|Iwhjut%=2bR*khA5M3;`7a27>)&g>$dW9SO4$Zk_Wek^Sn@=V`5_XY2GeOFxEzqX!;)=y zlWc*b!x|TfLoKfW?4_z5G8sy0e>&YmOQz@-yE0Akydar{1k+Z*)Z7xz-X|jnDUSpT z1W9J7$;YT5rtvTI%*WK7r`FzB+9NTWjUp4T^B()I#S$(Ldcu z;z2+gwJ@r!sj<+&fxN$};=X9?)Wufz9fFW|WH{b`a{b$Win|wgCpZ*}YoKU>P~6?2g%)>rcXz$%_nmWo zdi9^1wUVqi*?aB1XP$TDnVE=>%JQ!;USq()!M*zM{+${e90CXq4xSbb1@?=s3_}O( z3*JRdUK*}^m|_pMK(drll7fS)ip6|1Mux4?9p7uaz`qI z)8MGh%a=%Y`o;fc{f5#oRvbrgm}lb$PdOqFN3b{n_uC(Ak$FKhZ|QzCA`$cwy>pVn zle!jTt48~RWWGXssH=r9IW0LWJxibgKPDyCAzbIG>882zZE@kl@saM%)qU%3d2#u% zuv_!VL)N~YbSn6ntpC%%O2FvBRO5bEj($47oKx9LPD&em9pg#BiyqxV0@!(9B9-ll zREk7ZV(%(A72JwMLbofrzwWO$|7a)1LTngLmVcNrit`e9qL^DFP4bL55jGtxt~GbP z`?8sDL$#Ovu9szBFzp4=owRA9)~+ux!YNc^D^h~LqxyR zOR~nL_|ylF81CO*L-uQpbZC9Yal;beFXEaR<3k`a`VlW6KG)BJUXtES3E>k4rP5T{ zW}iYJ{!NZdvk#y5li6NrD#uSK&+B9fAQ?kzw^*D$j>h@w+oK-tMH%G>8tEn~Nqes@ zLIt1Tyjp<*U2dQdgMX(kKJ_yNQmdI!KbQR<0cA4CxFS5WcC$;8)Ub^^W*?E1pnJIM z^O+3BJM_KnBCLcRQN{?;b+wIwSF&RYm}<=b^KFnpxFbQKJ~S; zLY&0{7P$fD1D#24?@!$a+c=b6okG7;2dRnfn~-eikeo2@JbiB|JOxiIg~#W^C9y3D z_0#O|wS;a&(onR>2*2FYB7YBYy--K_<3g4A4$G61l(d*ZDTn>e0|>pj`Q0jcx%1*F+B?$ zgW~_9h4+EXv`eDNQs)_c^msCQ*1k=-%sQW?FbX2Bs(*W(Bx$cxxxV+PkmE!I$vORa zk7G6yoGjX?@KA|$6E`9LB~>7ZS)I z)mzAF%*Z)rP9MjmlHvkc82`Dv3Dy7o#k#5|=Dm2!NkWDKTmRf@Och;vs)3^?WEqHJ%(HfwR{k`CbWBZCMd7aN5?J2)m? zLtO_t2@|o@nDt*ipo2uq-MC_5yQ0|#<3wPojr7BG+kVn9g?ZdX`ck)3C|->2n@s=5 zxGUTEf^Q|<@3Sr-^d?^vqd5f#Wm(zZ=kit=5DT4Taw{$dWA7bL?V20)Gh-9&1s-}= zCZKA5rb>yYa-K)U>%W3iGdEKBB;8GT6F#iMj0p19z9|O%H{i<&+DHEd{&bCOihX%M zjn{Bq$`|Sn|4%i4Q5kg!(|pul&A&wi6=MPCs;T}%+x7oyTl4U|ar&2N|9@f6g#Q2I zEu8x?|M9BmW&9t(_y5N&Kox>m02YG~F4Nv9G9Kc((3GuXgg5{3%G)K>X*3a|5HdiJ zczLXekEg@J)%S`^jxvqOH#?6=dzmF7q(UqXj8w7;QHoSp9~b|x}a_eTJOo}mHycGNdsMCcmMk*tL^ox%9bBU%GobM zZ`WjVUESm&Sn>XE&r1$uc~iD&q?mF?WXgYL%bMnEvMu3`{3~24Iy5Jyq!b@KIwe~& z%7y^)K?7#eij1LgVic6{DCk8P5iAVf(9?|cR~gO@@Os18J@~Y4sQ>L}+>Mq>S2u0;!o~2F>xg+J=(8oO|G48W?@ndC z@(-&y`rQqyM2XLF2cog{-!N?4M}iwBRQV2 zI@$V|lf+aWLhb4{_LkS?1IM9|C&IA7EyMpRIh#HoGp^PbX_DQBepRZ_NDF*Oy0@e* zi37#~!DMdTXxjkRtYa>2?`4lapB?@?ygmh2$2%k|0gLc@<>}!MhPV*n?_~tSjTFZH zc|%!ur|^Mw`pwOP8lKct;iXi5G*5{iDiv5>s8X;`uw&^Yk5UF$RpL!^aodl9z2^o@ zw~o`=N`Ahur@+{j?z_3~6xnbu-U^?Q8WC6yxAd+ESaSA4P`s7WBq@`wK?fc&0Ir^J zI^Hp~GAES}GdnQ^DcOn~1T%A4+<0_pIplBsI$-P~@6qha zAMtWuOQwyh`|2AoOkrQ=_g#_SR-z6$>=on7`SYyPY~N?dqzzgW#n+=2$(O18XJ&*< z*HX7vTTBkwj|dCxWJb}y#ydJfv{5JBBfdMD5((FuAkx zxAy>^38r5Z$Lox#Yoi;nRSjY=PGk^3PmEoMpWmzoYlk16BhHTs+OlJA~p_ zx0z{K+B{3_3BZY``AsibA~PR0q5=K*?TMlue0c|%YX>}hMJR+5J$*VxFlJ@0zotZo z@7(%|&KnDFMsl}5J>el3Qcj7lD|O>EH(S9>@OxIF*#vX&FNppILnqIVH^H<8iwr1| zv)H-2=I~Pa*t808qg`ol>;RakAf%s28DloPym%>R?h4df%qrQAlv?NPm%5IN_nI?| zs)V0QnaQAz4_??G1r5K3`Pz43AyZK@e~2R#1LA3I>=ewXd(%y{@~p0>3pqH` zUi15s#)1HBnj2jk`cRpk;)gR3h7;6v;FXC94eQMN&dzc&6W>sE&qgg98CnKHaiQrC z7T+`MOfq3hZNSSxk+e9*!Kz+P!xAd2qTtcVT`+%f zO5-~1o8HL1w^@-ayManQjabM`JkOnWe=(2a0vFWdB%$$F9hJs1-HBPuc%uKtVg&HT zOD3ps(3?lSBQ6S&i64iu2j|LAgjflM7g4rm`YnEj06T#h1OtZk!(Aac=2s^ln;te# zb#~91y~Qrm@5@8{oLewHHpPA#H^v4DN_-8!%ndM^a)Fyysz+}9B5a9bBgP^!XxuMs zkTOZKv;)}M*ZnMC{#l(E$OH(TxohR+WJVI(qDXk#8{TJ3hs=obG40?38GQb8HZX{y zBqo^_ctqECxS`GFE-G|9P@zp;V5m#P6C({7^&$i^jD|cw+@uA&kL)O_hw&=+AMg;F zn+)vU+>jX%adW@}DRR!=&vlt^>{Vxfl3yV1(#$+;ls^65x8c-CVO4iw2!8lyi@T)Y zLxkTM?bES*k@i$|HQ{xWBdYkw{4SU+!92=Pc47ZsQoF0SA@k~tzN7DuiWC0Z2s-`N zZ*#puAuFQQXJjogYZ^Pv>MY4I5+wI`Ev87&XC)blmH% zX@kwQHY>#i((PQB-@Yp@6w^&Om-3*o5O4EZ$;i56BctHsuiH4piDkpp-TjETWk0S` zyh{41OjcR|Us5NbP;U^6w`e$vZQi9#iF9w(Tge8EV!mTj4bWwUR!cUXZD zN!oy4u?m`Xoi=L36!-<(zAta&WQ)HiF0AIi}&J?9W&zQEtcdVP)cSP5j4@B56r>sqD{-O1T zHIfsQtn~j_(G4YVjJCis0g{D*vZ-TFzj9B~Gl8*CU_K5SJ$oNFQpq(|vz3pkGHnS{ zElpfVAJsNM_95OV>N-NvrGo@KsH1E` zK>+cf0gb&{NW8mGqaSb8emwelQnb@{@`iTeE*vB7QEvIQ!z9lfUi3r(ojqIWaYg`@ zwl(zoP$;5P_`Ngw&Cvt=L`|CTfVe`1tp5-jSj12Y76~k5M5>R`wi6MG;mqv!0}BF< zh!!aSJpO)gC9hFyGx}*Eiqcg~o{tAHSwPcnVV`xrI!S!I3@+!;3Ek!x2<`4+TJ4Vg z*A2ClleRw{`gwvci3?gCgQa4VZmH;{jm z4H-(RkfZ7Q;N)-3UMrnTJ_R)Ruh-cACh?<`2UHKl4B~SZcmx^{{B}Zuay-5_D!#$J zweZNl2-oy_pwCjrGZm{&=ptOOg!{F<;b*0sV^Lw%LTwzDO0*+Q{fyv!q0-kGo~<0J zH$OwaD{%59A)r~HRlyx<^0bFjJ)$JT6kzy zkMX{Ey}zg>U8|UPNy{-(a1nx9Og0JiFmX1?N>3!0p>AI|=q~yV&V7MD{KWfg=ewAq zVA_p=GRH`VH=>jc`Hm5JP z4Ek>4zVC##mXuL3+s>)!;JqiuH)bFn^Z3-K+Vx*dM7d_KZM;~R#iK@?Xtb_qM{22g z-3mY-x0y8V)=>5JYfS)dA)S*d`7{`7vSSfWQ+N=%6<)@{Of=4RA@u3fb3_=h;~id0 zJYI^GQjBFz2oD&IZD3o<6dv;!&su6dqucoy_2Nz|?}uNjEa0~Fw~x*4@Z)=L4G>kH zu5wl;AT8F&5Zf&5njw>~XdIw#`i`NBrw3m!r@CY&@n~dg9mVVb!`i30W>Mh_UYR|t zBHaGy1Cq{%QeRA)X~`cP-C8NNU$N|<&*OO6f)~ zA9u@YTfg)qwX@HOx8g4QQxv%h1U*`qh2+tTMa2L?T;X-69geC+1{yN;;#EG=`Ev zl<(_&6i|vSPhcUHqA{mLuRT*;)&1Afe?C@DpXKPx!-R@ocL3(U){m3-F)%F77spga;Y^-knYDBD5&*}DJm(hhqsqu6gN68?WR>|_=j7T3)0%|D z%j;yA8EO{JyNyQ^)X!9=oBdf|BaTOS1j*bGYca-eHR+<&oX3T~A?ZwP7RD)i%@q*! z@V%Zx=T5x*$z`p~qp-(9KQsJVi4r^8f_h^5Q_}Z;(5MY?i#UAu=HGOyJp{gxb>;)o zw^{dTrpWJPH9MXoc7eInDj^j5nJV{L0Smfs1mcgVjh<2wjYoymz|QEt3Z2%mP(R-D zAtofK_-+RVw#dN7d+_i}L0=dBe%?9K*#xvs&9;q&+NXow<8VXYU<&AVGGjmkaMxiUtV{FYwyltB?< za$So(&CQV@-oLw@6~>+D{OL+zyI}1-6fz(I;;r6UivpMB>sd|*L*jW$epY{CaJl9D z_M$6{^57Y6>AdO8Appz_B=ZAcLhDPIL$4x;q1c@>+5YIJ@;Ef6hlr-rWuXT8Ox?4C z*^uixS zBkE6(Q26cjDUh@&bS85!*w~OQoUVQ8gc;9B%Gxg>Bc|ew8qGb%yrsy8kJzC^MW^A@4{` zjlZ|m$VU4QLZmyn$ua+y{_zo0NQ#$^MD|}!gveU^w&b<;74a&t4 z)365xV(cXP?Q^fxAeF{EI>Rx#dQ*?rxYzJA*yP@jV}+HP;^xFui(dn&I%7r971155 z?3T+nux4WxhzcA{#bjxU2Gnn3jW>;Dlk^h8@gr-yhSl```F+Ph;*^NmOK_IYcjE>}*v>szaQTOJC!hDATw#su&e> z1{RfREYXwJ?vg(Ga+BP*OJHnMYQ#R>b&B5pIKdF4PyJQ%myuD8>{O=;ECQLKekV}I z$Rt(DR9)}K;IBU_#`*s9Z{P=J4J>j%14dMTrGWnC4=hq?;RtD~)srN!@vz_EMRf3okc5GLRpZ_&%X?1* zphUGF250~cD0C2MbZ(isI+mrlMp$8bRFjKiQMdVz*tm>@hp*Wo&eNHJL_N)2d6du| zVgq?u8ka&YI{J%nZYW+JXt?P zG^9wnddS$Cnq2~s2M=1AahS~nnDUsGn%~azALt7w*z*~bByBa=A8>m<*VM{B%*Ld~ zrm4(J;(_#w_ui-gwC1}GZtx-fkj6fHhpFr!bD6O2+053A5;F$FYhAU{q}1^4m@8K* zEb1x>?6IesN|^MR5rD_Sld8AZAcT&i!Y+$3ur2YR7~6huNu~HBm>`O#(>36jKyV$f zRW6gPYy{#E+r;ZvpbwVT^>cK?@+lDd?xTWa5oTb~Q*8!3E<4i$i6C%n zV`UGC(rIB|rUH|ksb>LZ#_ftYogaVXu2T@@Rf?|x+|*e(gJ3i%COUA;4?)onjcx0z zxAWk)m->U7HE!yoo-mdcPACD=4`F;bmz0OHN>heT-kbs6S?L=*sCc+$?@7KJY9UV8 zwjU{Qu8$)0gP6QpfLcRvS*19AC-3dVpUU{-?C`obEj^f~4P(uwHDR3~u~g~!MXJE! z84^i*l0Nr4)s%bwycniw)g}l-5!v>V>eIe@YsD13RzU_GD=ZD~TXu6Zbc0uty zsR7&eEF_pBF;AE|BQpzY^6g7&e%vvdr`|qYUu0ca^oHiLUY(3{Pm&PUTV15FxqQ+l zcGoMM?m3NTU1MH34&U%5sK^eO9Ap@c?VN%v>g}ow%zEO%j?MVP>ypgO__$gaN?fjm zBOjAGIvgEk7V(WEE$LtK|3hO{)?gYEA~6NyLZQ*tX&1oxYkMXL75XZC_)^f%Zl__&2q67N&4hCeFH)e4XxeL3)FwJ~Ue0N7Z&@-m%UEboL zxThm|?YOxSF`YPNhHM#HbUpT%J^U2jeVF5&Ew8Rh z7mRhcHG{o23;0GamdZ!H_8D*2FT^(u^~uAD&7{7~^mS@BH|MK!TprIN3oFfI<#}H+ zzUpARs9S~a2x3zA) z=h!?O-rQ-<^yadpwdH>5V3=<E+nM@Vt@~$_ zfaj%(m029inhio~dA@*?8yh|p7zpC;+(kxKJvc}XuicW1C^O5@Bp9&&pF++*4iX_y z_TkXS?*t#6pHqJ%Yu*OeFEUppM_pK(%$5CbN-&eMO_lbV=!YE7Le^#Tep0&jNwW%J zUP~^QsIwlsNp&ZD+nj~=0bCRO-Lre-#*0x(r&EJU0PwWO>DA$hQFq2r9G@j7Frx*z zJ+p&|k$n~=v8JIbDRjYj|N9ZZZVVgJ?uk06m4L&Oy4_Q1E@FIro(b^y#Q|yf4!2#w z7JyvhE_1!VR0*)eVbiHWC9YV)RGdBvEjYODs=j(K@;U}<#WzrE*fWQGBF8kof+EIU zEWvAC=;M_x`QlDH2%$CU;z%vSHlFc=@oeJ;z_wokVFk}_*^99LP#8lOGp4J4sn6&k zYopTi#Oe2nE|$nma%cGy!^e-F2=3H^@IH~s-OIsDid%^^{Whg5Y_sr&33t>o+lDcL z74NYs1v^P;=(qx$PKR2{M2YLW&G2zR-M@O&@Nq;Lw?`qc|G4~#A~5j@${TU3q}ENe z%SHG#>_j^^5KV(WQ%1yI>3)ccyNxsb8A*Wu=G%uT!Htk#7`-33q6ru@zFa4G&dINv z`@?PgvY|M4kp!dszTjASIoBJwY4aT!xM>m1KKF}pVLo=cdrsK_wpVQV2y370IKnVN z2%SEp7x;6Qoj0xEtlI3H2TLFa5A3l)CM!0Bf6zcChksoCeWIs2)y*LI?HCHJp3!eq zjTQ9#xtB)PHXEp=LB!kL`;Yj)C9Yj2byOYyOYd$WT}iVyuU)vilSqBImsSi+`2a*@ z(2$BCd`^Rn_v#S*nQY*Vi{Rh#of!g%)g!6%I!{8rY-{to>p!rB%w&d>0VDh>xR~Am zBdVAto{spYBT%4p-ASjPKCZ)_Gy(&(exJAp#%Ahx!cBi?WPNvyKkL0|sjIz*+iZh7 zm@?39HCK)~f4=L6<;~~w`rdAGimQzDUNA(OR&mCCaH9UJ9BWv8NVq)J$NuS&BP4-a z@AjIwsbgBtH@cc6)aU*NtWvD^J%%rr{84`z22%W=xb+ye&!%i|V??dqWzIl02%!Ak zg*PuaqO+VTfkm=yR6#FbGkP--uRuKE9SE8i8}^?Yen&>%y;2T?q>^~!O zd?j(*pyKyx02x?D0q1dXUXpGpg&kg$+;(C)BY)X2Y{LRXH%zu{WuqSdlHKt7(MCPK z1iX}iv*r!fe_1l)W}8lW*?W4W;nN~fN3q2h!!@xj`R(Pw|FN+{@ z0O}eY5O1UN{Gx9(J5j(-@!OWh9rvdTmX?@hKk!kK!htLAx#q*H~j-M(x%L+gTVywY&nyUTKyeq1DkwBnc z=?#gH)iAk{PBCnV_0^)|ZGU{)eepow=^ij8EuDMQ{KuFTJuqRzWbB=1^RPUCmpj3# zM+5$>+FVpr0q!&23V78&D-e8c8KUM+*evlbz{U*Emg9FIKaOj#L)9z$6<40ts#8lX zo^No|ua2F^vS$SETeZpUeeLO@plL-;r?*EAk~e#=A+J{`Uq3xz-+m>4hCzEM{U1bl zUArfL*MGd2PCxxZ0HsJP;+!}Rn<+dh#1%cBVx?7oBD2bx?+#+&0)S5WV)`l-jmfP6drSS!eoD73fW5uQFhwOM`PmX2&N@)2c*SwiQ7qKE2PF%UnIaIoHxvo7uJ) zUjMwCxxDn=7V|T3$9JYrBb;}#>q4G=-JWh z(#MXuyA(RXF_+NQ$lgtzs-?_Fol+OBu>%K@FV{*7<1XjVNj7emMYWh+b&II(EyXca zpJGIJ_;PD0g+*`qqvkn|Rt#{&wMVWk{nAmyrnA@GmQ7?3i}Q&1)vpgW{stDNjYms*S}?EE zo4zGuVcj1xqJRwQ=U%BjcI-feT439tmfe>%q}(xjG?UnW?m*j)_C+-yO|oRaP&&~g ze?I3JP%hIIBhwKZB}4I_{*4S`ZC40i$bL-sZLAdhlNJY$O7`r7k|mI_k`#GaF6YTF*MGi=ybOSLPt(+0*XF6 z$XNFgRbG4~+YS#v?VdEL>M&(_{MHEzg+?P;j#k@;7Z%+%8+T58(NX-uADE0LTd2p6 z?vT=p6_0jq{4mm^2EtVfb8>n z2;`CA6kePxdH|mi%f7h#bmVrH)EG*dxte+V7$l0?`v(+8;w_u zDfWrYCybqlNcP_GIHCty2G8%`xIeCUopVzM#;&1S`8S+-dp|!&1f$FLt-2sseCIsK zxV)cr(u#SCxZtD$N=B!$`U!`)%_z>WP+Mjjilg-1;od~FvOgf)OxcvPecs_)tEF6q zRVYoqsn>tz=Y695nmM_Q?D|1*(koi1-N+iz*ZJit?s*6vqALf{8LOaeu`GYVg`S0T z_~BUekbKNuQH++er1NdbWpc2SzH$?}ZAQ_c{U$2bEKwBy{Fr;gLh#!S~VmD^(9^)0L@ zzZ%+B{uK^g-d%00TiN&}iwV!h7nqq_Zu;)$cI$wkNkDm|F?<{68?hq!@^o>vMDsO; zI@BH~Y*ko|W8NzoF&oiux%kx!HRa2j(`ot}5zvy3da!V&i}AiB8I)M2#eW%|6Wm19 z3uf+_bl&w;QF>3FjYA+39$qx|*`vn>bY1236%F()$D_^C*h<{T5*Ca7yweuIpz^ig zXG^-a(RR`ynCFb2$7lZrG}=$JQ%GYtm|c-ctdh4PtMAyJ1_MItSwd8-!*hg%!1meq zr@Btdc8Y6F5_5fMW2RLdD{V-K0`pryCL?-yz}=c z>@H88X?!XQ_kacuYRjv(&j^d0FyhO$)&a;Wsk- zUVsJJzJQ(P4@sI^e%Ya14&q35JL^@Ac9z8u`%SlGrlnTKta=$A1dyTa)&sevOk%pp z3mVqvc0t4LJ(>Da0FhUk1IeZTsa*3QQ!iDNpiw&C4=z#u z!+d?YW8;vbPk*sR(2odxCu(G2jSpYKnjJFCg9M~~ua7CvCbU9CVjf&X1~?2A$vulx zx<7)rzLtb)KQ@sf*cfu6+tw#8u^3Yfk8QU0;!p1hyi&>~#eabN&s zM7R8c54^I~enqw09KU?S4SeOfrto9M6tiumr+e}d-tpw?u zc8xRV>p`rbgpIruA_5<#=)N^x99d--hc0&!a2zgEUr7lA8r7fkC~tWPkHH~NRc z=|YzL%wAn-{ibmBj7#2=obAp1lIiKDY3$kcYdd|6vT}kaCzc(T_byA|zR%1yV2;6; za-?aZXPCp|eFNarhMAvkN)kQu0Snh4UFvU0HCyMDHmh~6Em#Zr_@rIuspvpAL)Lk! z;1M*!wf1qG`ZsWwL~{m!9}K)SA5E@?e)V6h4a87t`IDmq@!)z)lj_wBU9)R#!e*rF zliS;3n5xUOj}vXcyK=uF8@7ch3ouM@q64d(7t0GMRm>FxDk2JxM$5Yv0JfM5*%I64 zjeK@2Z0lEwz3pgxhM9IpKG8z9Hxl?brj3b z9VI!&Cw%iA+pH7ZD-$D9k*o$_RHZOG4I_nToA%Hp@){4QXv!V*rkVQr`p@fwiaU)P z+K0>;o%XjCo@o+*OmLUkgQd|WdEIL68WK0fQOGuqZ^|_1w`Iz#)cGOmS+sH=kKpl? ztC^4cD>f89xkE2YC@NQm@}@?QrCxDou?uTyU0%M+eKHd-?)M@{GcP5`i_b5MTUQ1z zNYWSVD`u$p0X^t&AjWcM%{#+duWUCbvBWwfHW)_S}P#)B4rtMM+NV*f_qB zx_SYSW?B>b8cL*tY3Z#h^_P(vPdc<)yb!5(?0%!J)GD7^BJi93xo!Wm)3 z*348MM|qs-(WXt_`PzoPIzFWzWLL3*U`{bQZuZ3Wnjwg z0J*nVRXlyr6%R0GJ*qLZeN?nWyucewS3S zSfC4UQ^uGiE%;|_>>`RD|Z(HvZgD0?`ZyUgXF;$J03I2ji=7t60?(myu zs@Kut_|WJwGO>f6J51hl#9l97J2ATxR-4gUU=Cd8GC%9yw{!g4zD_8DmurN9u zz)7-g%N#i6;|~jJ(qWyWKa#v`;G3BB@909Pr52C~+v&|+P`)sRf_M{Ai>|M6@OWu_kCO}Nir;1T z^>wAI-O{V_xk;KfM#a{Fw?7mN+msrnA_A3mRp25G<$ug4p8j$5wZy1VYX$E17v>yf zM30`#d5^l&!B`>k6<=4JiVh4a$~c=}y^|`(TU#B$a=`*>agkUcuhBbur}l-_wmW{; z4Ylmsnu`wi{ppvXR6lYd?gW=9EGl%9vk93h+H4i)xowgfhUp2$1 zo^ld9EiXPRK<$@ghoQh}f%9vm*5^Rt{F--qDDJU#$3>B|D64Zb zjI4!g%clvwD@<=Z-$8VQ2#+h~&@GPY{{t9&82fl}X+t+8K>E&}zSne;CFYVOYO$dD% zIFJ|lHJiBieEqtF6R<^uv-X@G`xEKw*W^>>r^2<1CM>9x?WwuArLFpOtUHDwGJ`>D zFRcXXp=GDz)h=f$G`>4%FboQ-K|)8caoT8zX)htOyB9o zh1S(~lVP_`S(C)@Urz0W4x96i0Y-#g{7ah>UU(nU`Hg=0AK5Gi)(4W&nwr&LB=TKm zQD0$=qH85NKflfhUIO^A(?#d0dxcZ)*3qCpuTo)+OWG($^Fv{7us-O-6);-cxEp>g zbv=Uk>jo}s^WlX;b^!%jhMrdJtSU<4O=9f_15>KoAr#WfBRNvu2nwFh1ZK*N>DB0A zxUVwrDW?G6*CqsdNG2LONqFizzoh9J(NAfR@?qA!k{ek@F&s92E#e_E)#&wu@^tEv z5&#B^M;XT97+UW1AQflE)x}wjaX{W*4w*n5QVBcr3xJ8K$6+i3?md%RK0lykvxpO_ ztT)|Qlf{03M5}4*?-_mo5le|dgBt2>+#!7sn#&AvcdHw!k+}5}{5p=hSvHcYf$0%X z{7<-Se@hr!YQ#meZZumvBucC$8ihuJqXvvxEHmRjsL1sxFUQtR!k;E0F?yT2y+LMw1S=R*mQwzmANZ%m zADLid@bOy4YI$D2vhI|SFzc4Zdh1ZA;vWTUw5Dsh|Mvp{u%m(C2U=c-IbmjDJuDLc zR>Hb^LjBPca)$Z@Cj%w6mj8Y>At=I$@h_R#;z~(K^)?KXgj2R`=sy|J{}ZB~6vBP3 z&OKZbVczgRb=ko9S=fk^*ITmftKA5f%aO2OVB9|6^-CPpC*B}I@wDR;GjPhC^V0LtU;gAWAo< z{5W79&#%AWkJ<9a{jC3ZltXWUUgP_U>$m&am)tYNo|~epZx2#K2P%SB6w81 zfS5-~~+KtEo{14ckneo1EvAQgH*e!XL9| zFx1CvcLk*2Ci!>uYdEkLq8GxA=fC;~j3VyEU0YQe6`1-tq)VOc%C9u#8h zn*2_6dQ)N;Hs%Gf(?X?Q@p@JJ+Azn0T;n=SO7M!xPjo_zI&j0aDvkxm*)s=;HG=ne$2 z>1sW51EEe`n?l&w{6h=*J$?p&D>(<`q*7-DL}gFh0JjIse2H^pc^telW0Q#4ZNJe2 zE2)96G4`RyIFBe)lR~Q>*?YL0 zSG%1tQ#Pw)&x6^DB&x;oJ0Pfs6)70*KUFGLdaf zKtYBgGIvup`ABW~{Nt z;s7hM_tdi>B%#r?-J~I=e?bF2l{`5MzD?F1R_QSQ(k11pH!IK5JV^0q<+l`yc zb9kI?P#(6<(;dNvvQ>QU7xWrg?#|Wc3j_oZlG$h9tu0hpB$IN>lRYIWcEuOhIL>O^E;)k&ehP2ivC5y-;OwaB?^_tRXKz+O) z_84`Ei#X(zLNu`;lm)>gnm!<*JkX-WRWTAib)`$=b2md9s$3^WfoK3gHe{~IhWT;@ zGx%IPO`IngCr?NJm1YNkXc`Zqb2%q5n;M$bUdh&Z0;CnKOEAJ%+YeL-W*J0#in)K1 zV_J1wtIn*`+et`D#y6iH&uz}IWPjPK)}yFogY@5Fq>DxZ+%hZ*)6cg$=KboqPYq^~ z_xZi?_9M-^9>P24nKzJAH+7W6cEZSzpl?h4bvP+E&_T5szOc>_Lro=-mV(;&Y6IYL zDdF37K~dvKbFE!I+z23HXaxV>bajS}P(LPJ#tz=msUa(m8%q1|Eqe$RH6nS^;3VTN z=!E;n*)zt0XNQU2AK;jc6WbN-p>`TeNczAekoYRt!V5qF)b6WtYS};2!knDtQ+u)5 z?mzL-#{r?Ku>>2{=n=~`f0c&&pjs{W&Hxa)noQS|2qSqcLx1|HAwbk=Oe^y{2}6a! zrN8<2U~>5Y&p1Z?tgDD<+4Wb(EbAw7?1KBHm{dTP0pIX#nQw%B+=g9TkeR!e+cw8n zqO<3Cb{21c5BQ!TgX%ZU#Z|-LHj%ZNMHywwHB;_R+gNWNH0I64IU#B(clBMTOfRwQ z7kH3js$$Hx>6*dGAZ7pHhkJ*Hd)6`>rw_wA^NGd^@Ki(9dV-MR@9j+HR}ND)+?tz< zsW!`9k%HJmsI|J+&V+sy@`0E)&gyBfd6R(?46%D#8(68OwSknep20|Vz#)_=yF?#5 zjU@>wx}hPg_F;06upcb2kByLUQ!_}25TBgeBTivNhz(_%1K0JR3Sy@=R+#m^MTGL- zRD9ZxN{T!F{usnkKy5oC8Om}nt7kclv4&csbUvQr7Zq6a`&?x9txhP9LD)(Qx0x&VhWIW5&GrOX(pbPjp=nR#>>MPCQ>zPLjB1-&K7L zq1%ET&+&_l=q z3k&~gqvU|FA!JY;ViPx_c#!aJ%VpwI9RE$Q{E8kzxSN}L*8tFt@>Zo$V;HF4lXCC; zkl&q|T@Es~{-E;1>xva%k+(6j6bze>=<@VrV`ztE_y>1io;i=rSWbWCTQuON^ZGY2 znL}-7O#-&O;r*TqPg5oSq!pFysu@bHS$lSbWu$qdjTaK>?0o4DceYS~ryc};QZ?>_ z2l3k$X|u8UGL#4tdg#9a%T+qD>-HxYg)M-W_Y@J*oFH#tfkB-&FEEvZoahHfL|I`X($q2ZtVfYPG?zx>s)C^q>; z$dto2g3rJVLRam-9PqmdX$KwD8q;B8P8sgVc)wl^Mu1;vS!RwlT{LkGSo9;`kAdwnGcW~Urg0Ofxo%)DZVxEN++*j;1w`iG0 zeWd=W-pPtviRj1e-Pmydm5la(3+$&llLA?^okTRA4H9PtcEP_plZvaOO&MIR2#kx0 zD2OBgpPInTU;kUQo+avo7n7p?EW)!@G_msXUi?s5X3hBfo5Houc%I9(ufm>}Y!t$B zu8%%~QE#8AVv>BRJ+nPoVBx*s&afa}m2=#e-UtfTck=IBDPCE%(vJ&9s_5sqKo@7m z#WG9($)f(v8hX9>st*6>X@tLd({}OyLETwKMcKFQp6-%Hx};$wlx_rJ1Z3!LkZwV` z1*BwvK~O>%>Fy4Z7)nA~=}@|xeWCBY-*}$;e%JoEzqpob&8!(1=DLpaIDg0S_nRue zL?fkXp{&s+%NYB|Nl3q%lOirTBg`c5TFA~(>gs=?kb1h-+S@s(O^-rC@{f=EQp)btWz)?H%&9Jt=Vv+xr8a{LPpZ2JAj)`9#nwrs9|>&QmHF^Z-yMYK&qd z6CYQgVypHK6UFlvZ?y)nG1FYY_%&w0KTrf*OaFC*1O_OJ4=H$ucBY(_fA}k>zj7r4 zz!0c!YT6Hr{5@brJN%)!4F1Yb3^+>(n{ra~=iYLInA0f~LAPw5&?@zXA~_<@N>ksP zVA4+}+=h0gXU#q_{v_b=qn!Hn16I9~ET7Lw-f7GCRVuThR+>hho%j$>5MRS*MfP~@ zC4a{9F6GU65rj24*%DUagrY-yBfbyz_|Z1Eo}uw2I}EGWr(do#PEv%<?O2>KpLRGJJq{6XWO#GDRQydn@5c589bZSS)142@7mB zg=9r_q7#`q)>?)Xo~~r_qM$(iQ;un&H%c`reU@cY1)9o0Z{&;kxHx|EsgsnQcF8CE zsHbM_!&lY_mI*InSXLObUhFw3Y{nfiXC236Ly!l9edEX|Ry84bK zPZMQ*SOc-XLDFtFwHqG6qjd6DN@?$`EwJ;7M$(R$53leMMxfH9pKd|46j z!(5R5ww%Ce^$F#nb{GvV9OMh-wEF1M7Kfui2l9L|i{ver!yCpghN}?ww1}{4{iFZ{ z}!o|H~VHly3 znjMHmV=XpYz90X-^f30j4Pa=XL(tO-wn8O@NN~LMSjLP=!zmQkUPCv{a2B&M=2-4C zk0>!%q#==SiD^Oj@P2txv6s0KR+SQ#_0%{}1f<2I(9|-HSjEO%$cHfZ2;ZBUavh*| zXgc8Dac&&F&$~H6jsH{DWCReV1!-uIUe#x0Fj;Rq_@i>95G=J4R+wykA-yfk zoXFA&ZX3~8G2SDfzcp<&X_S!+ukK=%!ccvbHEmU!98Y*2nSbQ75k`LzJTn!5!G$n@ zm1M7;xWV400iniq{)F$!4|!Py!ncc>QtC6>^Qg

*M!9AS>kZQ4If1Bta%y<-6f4@8<6AFQpV<8DvX1T-KUn6w(-TfD4BPa8F6sySDO$ zBXy?SRIfjA?|4HECvA5{A`TTy)Z+({?Rxp`aP490JW+<}tqCCf`pO5^1#89G*x;H~ zsmL`5KW9s2J#K>om5tTaE?@>I#AbKpTQ9gAm{W&5uOQm4pQ^o7*iu@Ky`-G%Bt_7O zE|<=9QH@kKHqFtM)k$5Sggyp$U!aJF1Y4&KM?64dsz8gFy;$X=5@sLRTYc16!i_!y zQiz{y$=C^@Vgo<*C+cz+OP$$kc6h<>14xg?7D4T2)@$-UWr0CH= zkT&rf_G^n4t{9e|Xu6EK3&eD@b!FjEvgsbQkAadNv&@4Ukp;y7WTOC^=$aTL$Y-d1 z`F8~HaTI+mhKke6{~TsUBZMmnc68QG9_)Xstps`n?#V4Qrd-NF^?;vM0k7(Lsyn@4 zDj);V&5KoGvg}LCHGwD6Ash{oxhN!^6kL4p5U(Emqr7T93`7o3%Xtl?jqyL@-(^rjl~>&mRShO=?>DSa2b+HFaU*;ATG zTHkoNoA2gR4P|Iv+kC+xoh0R#?buTg5B?W^Igm;z#7UA*>c7&^=jEfocGph8qxFRKRHp(XOP zk--I{01^M&XdhYsExA9vz1Fif?bptU>uT5B1oXJLnVoey(cK7{{A=ZQg(kd$=R*j# zucs$iVVN#+1V2+Iu4+e~pxLB_QxGAVl|r4Y7v(?Z82Vjk`xfmec=2+paqv3p0bg6~ zAom$Bi13{S=+}h(8LTYyjZaUBCqBBb*ATu zE6}S_)H2dqz}LQe<~o@-7*@j%9=s*JIar@G>m(=V68UJ*p4?nAk~upR`mwoe^s*c& zVxq+S6jB9vqD^SCk&~gE8b9*iNj&P(a|53kq#9V2z(2Sjb4*6c^0lvwHS-3tjF)sG zG693z+xOgXi6*B^X<7GRN6Y2XhI*kIzSCJ5=R&qSIsG8;sBNLr2Hx&ku+`a}JhI3c z;9nO=U_adNSqrtubJu0^Y>MLv(dAH{B)_;xTcN^(_M0sD@!$r+>Vq0{ai$x zxxSQ6JzmeV=scH^R6XVz3TIC(QwR)r(6;L-0nFoQFF0zJXyE#?%KdU9J|+6bU^?i8 z90vR<;r3AfDd#(UzMIxN^ZF94lFpx+oThtV+Bb(q*j z5jm1eaWL`c@~ixEqh4BEp#ggVx_`1D8C)D#V=}aW7X0fALHS1g%mjnmDK= zj*WEDkAk2z2WNYMH#@^ab+;p`lvkgP_C`O?v#8c~Bv`M=6{_0XJ|B(%K(=H;=1eCX z|AU+;3KYogjv}(MNxbaBt2Ehr==8|4q;V{|OZ@>{^I8v18Ij^bC{GeQk>Y3#s2iQ@ zZE4p<;*2b^7p)j=zdxYe8S5lP4!K0bVxYk?HO+Fl#9gVUu!NZ>j9kr*`TbCt;yX9g zj;pCznZMlhfevU6se3OxR_oj|y7@7D{t>f-6j>e1OI7+aKkjQ5DQ;?lYc%O>K8Sq! z=mj$CC0Wmhod}V#!^2m=$AS-Dm1~1~z(m9(ykMedLW0EAuk~_CL`G_^>&nw7na>2z z1XFABs8|6$b`%ZVx0+4OPr%$Z)FB5*l;C1<^Vsff?*8WtRI*{ zw&(aO_bJ1kI&M_hn2I=)ECodnO**WQZt6c2i`ce%y6&;V2EQVxRVL}wth!f@Tle@h zRPaMyr&a6m@`uoOWkV5N&5pA1K9{EXp2O->sz7^4&AnMV&_wxTZYIL{8WYc;#}Bq@ zQ_xa8RDuJKrYw)CVZzdhm*PCX!fM*=4K0)sD;yb+ zbueLjd46+_aoiKkqpmwAqKj_xM{$I1Fu>*JT9n`esBx-s1Ps?8r{52lR>*5Yee;T; z8L@7igj{3oq`~momb_w;%|kaFE0x7Tgd>AH9knaXN|!z^Bvg7)b<|#;4$**pttGV9yMEj5Ik?+L45KTRQ#5 zQOlb(QCno8C575;@dlH@vpI%_s7A)srOz6jHfJb{U!*XQfLGndw6bt)n(^8c7g|18 z<;YKy5QDdkPDBK;aKQJ>ynf!txjpU7@@q0|A1{2!$De6TAJx6Q0JHr77Uy@_!NXuHm8A78r}9W{Of z3O7w!Lb|60lRZ;PTzqP@_7SOzmi2oQaeFQb6)GDpoLz8#d1aYw$`w`T{x%^a0y%Wl zr@K``_43`ey+-^q!Q`H2Ivj6P&O`9-80+lwkq3Dm2tN$?60<9rjl=1g?VFE%}=Rqa+p~$a8#Wr*aZYA|9 z-V_=lCi=1)p&xfEiraugm_H1fgR@m{ntCZuRF^uA}dPN|wytaR$^3Pu0K zX=%8(Au{Jeb0KPVsFTqFS5YdaraZ_=VnS_g_qjs zaTwdQs3>^6+~q;OpbB~qV$gacf_93>o3%;~{NW-8u6t&W2#wfX@m`E{6ZR@(UUzW@ znk{{hz7<5V{&*8}nVr_WqNtf@ksSDiAr~pB(Xgjdb7xfwiUM)_){oj=4GM36P#XHH z@sWU;kua=39cj?675Rr#JTg4@G+CChhd8qm?b8cqnFKt@mIUWR4A!FXh0lo9OlkE1 zWf?sU?rvck5Nj7YO3+!7pU%l=boq;uyikrX%wrpDfB!V};I z+F*5bf8)V;FTfG@;0);Lzqzl5sT81`nZ_?QL#5~SK*7YsRU`F$>NkUQb;kJi&gL!P zS0;qO<~uXhu~fBj$c}!)8~p|KwDya5-lvd#iojJiR7h?AqqyXe?AUjXh1^=a;O=;O zrvFs9$6M!_3rc;bF4egQKhijy-=E(ovWxrO$BEB~Kvf)){`wpqsyI9vxtqQGqsvTl z|MclH$$LN$M;!9@M7>T9l^0Z9(Pb2sU0Su3US@sX;|d@rTdgTOo0ct^U4OmI-w`Xa z0y?W*5%X+W3me_}HH8Zq|2(_oXf?PLudrNQ{3?5Hf4NBi&zV;8vZ2sr)=v!Nv#u*l zKlsU6RDO>+|HixpoZG|dKiOE7-Jq{^!1P@xwZ1mart!XKU9;w|T|l^`X}h|8XZmIXF0Kn~7ZKs$`k92?AzvYHnb9Y zBsD-dk=3m^k0+31A^ky)(Rbjjb5RXag9?BeMHIXlMTR=usHwg%RmPwT97KnV)-Zh5 za+E~pjBv&SpPTorcX9>hV?samdabcwX#dFkf%4v|iSSi1Ar|zbU+D89k?FTe=xv(UzkML{|3*tAe2KRQ$+iQ$TR5aKAP>*FZ&=qO@^?O zlG?I-UE#V%BwL6NNU^l0@(6CCc312X(=7FEv!(dRJPQ|Y^EP6j(6;Mbsb($tO-UY( zR7&E*W7{;&n&fAaL-@m$K*IJQyx=XjEtBm&+^Ky3N_6A=13uGw-{a2_OdsWcN|rLt zSW8C)VKs<*ki1@Vf{CiNy1VG8E=v+LQlqq_u*LBKpXmLdt{ZwRFxQ~ht;~7C# z$HP8<;Y$VgxLLh58bxb0eZ?w!St?V`Gbf zS0^d1Vrn+4NtS}eZZ(CU1?&-&$o5#l82nDbQK*6k1AQ>Rh+QO_-d8jld0U!Gab(DB z6GKnFUW=A;lXgP!6$qND@h|EDOI};wpHfmsltAWLdWX68j-)*MXh1@L1br@{3Q>oo zWsmBvCD17)IL@u5!b+{dQNa*O^q7o)s+AH|sORDOx09f_QQ?f`H?x#F(r*WJB7S08 z`x#dqKg>FNf(d^Md_6J2Kx4;oFn!1y2VdNJ;;=I&$@sPqyU`=Plg2?_yZmnC5AOe|%tt|LX&{_vzCW zOJrC*Vo^B~PQYI*s6m5Z=^UdXLwbcZtZ_Ovcq_(j_-^%%Mcq4~__q9EU$@`j*>8?a z+caF33xe|e5|%JsIMG?pyU6RdS0ARV7C3CUgb)hR#3ZMgPm@~}3?oBS**!XS#CwF_ zlib(2SXD(x%}}aW>3Phts8Syad-BF#Y)69lvY2)A?z~k%#%Z)afl+`N9di3e7BDC= zp2?GkZKRxUEf=>Yzab(S6kZhJaByYb<`@24Z>s#tMj4s5K<%q@u=Sji$<$;)23H6@ z=giVMoGYmh;nxXlp4x8Dd3tBap%fz)p&SU5)&XGbBu021`Do@5SmtgDKrT*mUW))8 zIhcD@#Ox=pJ@(hMdqqg}PeM4*C36W)%{g^JAhY3J0|w@gZUPYR_lZByi=H55TTKmU z=5&%8e@YJSdYD&;@KbJifB}8gYGM?ujxpW|xW<$)CgJ4g=b*eJX5<7qxV#T?@vn2b zUxe9;#vvrybapL5#rN;vE`1DXLAMngt$p7M0Q$}C1ZwJQqb2vyT~ zrh{uBrNZ>w^{%v#7X-i>c=s<;%b#SZ-Z>W^Xfs^|4_=!2T@Z zcaT;Bc0o2C`=D~>Ne&U72f~P)n@HW3iC?1Lk*}d|go&?>5jusicm?aUFc2CD!b5gE z63uN5H}lami-w^;S-jvG%4v7*Ci9@xNb#c+G)cCC|}2x&oN0h@|K zS3z2pqH~5sPl|KQJJ1#9GKk~{6Fon!1#!$gc@y4gY2F-EYPpvC>*OGme{$-3VxI=0 zpFTwkmV?BaU*@65Ay^iZGs19DwIIViJ&(3)bud7BKO}TQ`xfPu7Ru9-;sNN@5Y$^Q z!z&jB``z5*z_#dpPO1P5ZNm8aHU8$`<@~$sC^75VE zQ`xBKmkMU!S!L9%Xkf*u64$YVF+ItCbGms|)arJcjR}6vK)k%-7J$VCQjjY5F<{Fg zDuU^n>T#UE>%NqIB;Tt?HP|FCOh5&Hy;xD)a><6V^!UrfRuZg9qXUYU@V4*a;#L>1 zlQUZ5Dq%fB?S*~%BmPo`{!qPjf7QXI{dP_V#zG{^FyP8I|GF6le9E;KGMSSxQpO=m zkoAIpk|%#u*uU96zpL}2aNrR~4`&OA))(;|vPZ-c8~mM?p87+>dT$Op#|IR7TzBO9-+qvQG#VYvjiOx+mU9Fn|o(* zwueI>%dIA#|K%ZzozO*IwRD%en_zFvn4kbZkg3|&r@ubkla~U(?!R%R<=+AR@am_3 z={fm>Hss#vvVi~Br~Wp(NcS`G;a%W@ic~Wg{TwD`zBDiIJetLn) zpSA9i;T995pD|?iTax5;AvAF-Ga2;Xs~c)?(I`~esxS^~rh9qdQ;gynNl|{_o}TUd z>|q;YzQ#)yhnKW5<~=4i-!vKn?SA5ZdnlW;vlU^iKKA^9<_Ae`n%UPTDni~r(^1WS z_Elx67=J!kBBSp#5nRoWeiN>=Ox)=;@!PsOlM(#>F~UEK}cY9|t=4UcD`8<5#wPcGl+ef=c7IV2+kJzQ_O* zr?s?#*pWOhPus`1W`#&!Y8vQd>#8Hx;-3Qi6WlGb-M`Hr#Z&Am6C+cP*ZV9#{s+c)A`9nT_>CUklt~Wxb|w-=8=c3RvsN5f^8NJ z{_1)bH;X0F!WdkDekv;I0UPRAq1e7@{4q;VS&Nc0e7m+y=z@{~Zn&B*~QM4VnzN`P(?$PjoPGTnJ1{eW< zv)ju33{zXSq)$IkC{m&Z^V8l6iZc!A;s-HHv07bw39 z)gR$Tgf1jo6zXqpoX-$e{w$P?nhf!uSa6gRWM=J(?$!M&%74zY@lZ9Z!VD}K*>pQ< zD79I1Pb|}K8w1`M(=o1k#$SRQ$x$&-N7Wb}T+s939nvq9w`nt?AX*GP!pyskhCV0v zJ#}Sjhrn5`cS*<`D6P4V$P-*yE7c5{PHptRVcQ2z)_+N0vid~)vKXdS$m*Zob3OAd zQW>zH%LQ$p@bQ2bF;m#3nzHGulPfW>To3AQ4i3x~r#!H6h=}dN`<8^WI24d@ICd|1 z=Xs6yMXFQddK}_1x(o7k+2`xeyscFh6cn=d)7C&svxJVY1In@uE+=!t2(}yuWQ=vn z^-I;uIg#bv@$uA6S2Eg9oZvy#NKzC$4jMXH$F&B6k*5IAYfF{A3O1)^&Z8Z9BJ2lhD}(^(Szc9Bn1vDnkxc z-u;!zGP7J8Tu`0^s&PP?b4T=m6nt#@@gZE7;}-iyEw}@O17g71|JmKz-!7+}v~z5k z7~HpD$wT;EVEi~cyS!x_(IvgmY59aH6yK;zvv0%i3Oeh6jJa0P2ibR2{N+_(F5@vX^>!!go(4~lU0 z#LkW`PBxdUa1IJX;l(xA8wQ3YDqb~88E=j660+*PuRA00or{Meh$X-cUl>ePMYJW1 zB``VaD#3xAu_&LXCz;q*@yiqkp#jB{&FCTx32WAb2PL;$rLqGZ^j19EcNc(B zsQoGMS@SuW&#}SGd!k>W^F>xaDmGf}0?c)$%Vgl!AvG0eZAYX*qsNzBAIdTBnR0TG z_X|ik6ns?q9&YcvgZ9Jj_8lJ{n9MQqJK+_Xq^^0&dGRAa^7)dnt5^cpjh$gr-b0Z* z>m+f+3k<-TWbqZrJwou9w>A6ezFP%_MYr`*aWJhPyq{8gV@W0PY~hvNRUyz{3Ud%HsA$8 zm4^8_p-i4ACVnFi&!?}FGgu=B#%o--;J%(pH=nd&SECvC@>LaAJT8la#v6s)ITYdZ zD-8J2%qgq*gqoeZrzrKWLHoy9CyAK)dNu zmTt<4kCm-)%(%Ldtzh$`EqjPx2t?!%UTm+VUO9u6V@A((P`FEHYU>J~Dl71@xf;4% zf-FeWIEMERYF$Z|31FR5Rs31n7+TR|Sgvdr z#H$BVsJ&NTy0|=<+}DH6TB|nK`JZ$tz;T@|GOj)MdrkC}kn*_`&wes7e1(wZDV9CA z@myf1ae(CK;%A-##?*dj9_p`&mRHO{O!O2Q0s2VsR;$V9PS4L=IUH4xA&XyPfbSAp z1#1`7y&*0*{{E+#`lPC&cdQONzITzRhyr%#2w@h#($EvuP~sPNwqHu-zsqxlCqeyX zH&cFv6#s7`W!`^*CGsPrJl^&;drUJtGAb@3BIlz`J^`HsZe-j&+FLpTiaD`;ihE}XxgL{t7RH8`zZygCJ>Nn4FC}x`S7TV6U~Yk?qeDJFu7%wD`Pgr@@XZ^UvBuosBK=;>&u* zFh>o}t^hmfi8TnrTwZro}%=n~}|Q@a=!?Ef)4Ao*LC=2!jz z!(VzSWhb$vgj(Jv-XCJf&>-C;_YWUx{-H)^w=Rh0Mz6)Ry`6YGDsK!x$HI=jPGnTL zi#=L+B==)}rItZ>2$M0o?*1Rv&`B~pN-K|%Wt$h8pJbpKc7&f8*W*x`d2AwZu zYqVC3&F6IOQzC1wQk1L3Z)+!44t`L|@i2Z-Mq+a>tsf%W8zOT;di9aON7WUk5nEk0hFL4Z$?h_rj6vl& zn%jEPv2kHeUsnXel51D4xzycPrQN7GXRDPfp7YZcN7k=};gAaTdM(@WNG}qgJWI2% z8mUL8p~mbsgD|B+Mhka z$bAa@$KU}TMYFprd3Ea&Epf@)MORO6s_iah!xfky|98>mgpkw?xv}rgQr-Yg_N1Uv z!xm7gk^|{3oOtu&_>YJwH_ty@F9!TZQS;@81m~g1R48yL9yEv&#o$@tPc!As$c!YL zq^W)GP|pA(HSb|jZ4@5yTC>N!Lw23BMe9V92C_d_!Mabixokg7C&eYF_{N5xF?adU zXpbDWVy0b*P74L#3GJN~@@w{6+-o(RyfuE{%vY*+(lV;J$zxdF4Ge7HPzt9&oC_XE z>bbU>H@=gzC_%~rfC>C)B{*$sL-BG-Lv^&1MD211`miPmfawAQyiN5ExvQkihqb)i zu@FZEt(<=gL%p(`Cl}{zDZ*+{DZQQdF!7Sp;()3Z)rH}gt5R|$BQ84ES1Vb_mT<+RUf){-4miR zBm!%o*C(7vh;kAEIHEe-bx}x)8;O$!oYHxk{m@7bVj!S!rfK{-f)K`D?{LyY@V))&0R=wNDnSVe-P8my499be`vT`HcF zwy&UgL{Mgs0)p_MCULvn22Oq&8M)(2M|VMj*<&CX!BILe?sRrBLA}t=_N9x{&&JN^ z!FYDoxqP zUzrw|8|q}Z{>>}or}ukcJLrBsioq|=d(W&NTh>;UvU7k0mCf+GN}i_|RZ$&31k0dg zB2N-n~;;MpIT17RpiNe5}81L221%VL`>F=6CNxNDdK*l_@7Op3F~P>NwCZ zOP+0D(@`*~&tV2UuFFD$%QUSF4@Blv04{|DAzHX~A6BS8QPNB~XG!Y6N7L=A5rqZW zd07DS0cVjPi_@QGq@cjul-{ z6c57>R3F_!ITYdX?}7`O?j&+_+nRwS(H+-)EA->w`?F0k2-n4UsDY&Ji)*1D?{1D4 z`QE#okc4;F4t%P83m=&s2b@GSVpx)R<(rWa()ua`pJcp}dS$Ji6 z+Pgg80({QP1rcU~q06uC89btXXSpHmzN4eb4fQ|Q&>7BlNX{s!)S_CBw&jJnwMp}) zd!_Bhf-i28-J5qL%XdY|u7@-qS%xXMeU5~;u%n?3@@L}T$HGKWaxr)%g`NGHFSZY1 zh@~y$QO?2;{Q+{yn@j>{k?yikutvBJo}3FZFvltc7l14(N=Xo__TAgBLE)>7Z7g6n z*vwnZ_q)jgkY`z6M|L%V_B<>*lWCfWQl5zBr_Sw}K)BI4dm8rMPt)}%E2@)J_^q*Y zPusVmSR$=Tq8a!rzy}>ir-Tz&b}ScQ-=H);__e}sr2X{O3!1_A1#a!o11HjN`xGPh z3X9)Ras`@wL7&UkHkOoESDt9jca$OGsQMmxz`rR>)=X46p4SQhyswwv4sr>w(y|#a z0n7R{HtQ2&y4pHp<;j*HRq!Bq8qRzE5ReQ`u)Ac$IGBCDXng=s-kqNLD<&ViE zSYbWf;=|E`F|7%a6UgW64t%fzQk=9_bB(OvnJ4W~TfQ&PpWsqB?4v?38DI7x{TjY5 z5~9M5*?V8`{)fz@Xs3=)R+rzt803ly+j=2M@i0BbSO*2kAq! zQv35M1uCJivYDb2fvN;(ApcNEk`WgeU&!BGj zruS56U~f}b!iZ8>U`!Fy?AkgCsKx}aoc8X^LJj(iL4fP-!QHT`lJ`bR?kCOP@2h4= zp(J-QUdDkB#Cu&h$wU;mSpSQ@vphtH=Wda3_7|&*>;99{o76&actl*D?E;Zs!0;sc zmR(B|a*XE2zCqM0!?@)z7tGWyXNp;qhBTaF)Wd${z#~jB`;uf^M0Md4EEyX= ze;)I7v}*hJ97y@@0>_k^DVCP^nIMg@+x=*D#o>0D9Dt7P%wD!xblKMv*0_?O;gQ5t zR2k~1AX#5H>pN8Ui+&1=va*)CJ>dGG%zJnC%}@v>pG@qZ5{y3yrr({=j{oXIhpm;& zqZ{D%)m@NCDU@VLtkl9; z8A40{p`H4VYx_4q^Pd`v(dI7~Knn&v2?BhU8fPx{dtWcQXwnQj?=9JYR0#5Y?0ZO_!|ntFFoG?;8SDv4#<@}{SD;K_zTBFh8l{Q z4Z9I?!{FfMcWw@WSxi8~gSa?nDXWKoU*pg`+`Bt+kK#LgW-Q4xUL){+N{b1t)6chT z$v&w%w%pL-h`7K-i>3zZV$T}0$qkw&apTL2)!K=L0_XAJKIQoRv9WBjwl(z)H*fbu zqnCiY+}w;aPU|FW&N|yYhGu}>1GKakD3pv@g&O4#S0m1cC${` zA8{^v6{LEj78hLLt)BP8Bk$|}>c}o)P+XwqUgLNv!L~Icn5@12(WKl)%D0iLhuH9y z0nHtv`rW$W_SeslgA2Apvt4H$gnR4kjV^( z%GeR7#fC?u-lMcjA~+2X66%(jRpAvn%wcy6jI(vv4YI-QKm zkZ=lo+Zt`t0aF>iaxfA*;=?1RjR~1D@jUmd&Ir<@u6QP2TF{;3jc;Rmj!7ufbh892 zR1O62`1N(X4uII*tf{59G*Z!;!Q2(divE<&`<4>C9mY3y-{p>uP0u!K~S)|89$vqEnNEK)wDwQ+o`}=^`7*N=+L0`|f$6$5wo8w0vJK8u~uE zsDk$wB?g@6m44UfX|MFFug)@rI5Nff)+OxHX5>$YI2A^K4u^FRYiRPg)F$6E6oIhf zCef0L(kdpu*PzufdTy1(AUQ|5pw&C?*kIk&con!?4E^$YiYlMKM3yP%i``CCKd1Nj zQLA@Efv?Yl&GB&$4wdk7O_f1aM)o|ccBbPNlTWj-WQ{6KzQbisGWctdST`e+W9`h8 z%1lfSk!4pr`xszfKEuc1siys16VZ>7-J8bIBuN9FD9^?k?n8bfn}%)yt{!N7Flxix zGli1x-MGmV)j3{?>@v(l-$lCmp5rmNWD61>>~_oxk^=35E8e`t2Q1=eh9E6LHq_FN z3b86}*!!GQ2?IL3pZxTl5H4&uQ7N60H0-U+)9MPw*ERRjST6X^&AIsyo5-W z{8=sH>~U2Bn$7qP5x$&g@U2y@Q)0Trw`t~%dSY{>3IbA`ras96K%R~?&Yzsl;jq}c zAr4o-h7jH*5N9PVt~zAB`6Uarra#^lEl+Gyx9~L>3VlJm?Z94%*kYCSrBd$$@8?P2 z=m@HG`@{7f9nBJ&rJNsRn%)mRZcQ2RsuuqRE(RH+oR5s3q4#BR_ z51y80nKAR>ar~wNxS#h^olH67pmoq7&6-c9?aDwU=ArWA;UVhKm(!WZlA+XPr*Up%TjOnWlzhxY_~b2}7IYr@v=yGj-T@ADuJ|SHyxD zHBzTRU7;JL#oC^`er4U89*^<)D>Y@qmn^V@gXiMi_4u-RK|u2-VfIt*^=X%gUDLdp zuhOu}s~+I~Knm+j9bQcxdVgs=Y6LX_ID^pb`v7K%BlGBr5=ejVdyIvTeP!M4ta8GLYyw8K32IksEv6aLJ}(Nu06YUWQ5B!QN)S-ujeyvY2x2Vvyd#99=}2ngSe6bl8x3IPU-8G1Mm+;KGFjgv4lN0W;b| z#eXJ7cJC|q;q-|7LS7h90M^TV#lqv8i@A& z^pU>>xMa=C*04|d&mz%3NW~^_r*RM`x$`q99HKOmq~zxgkQ!Y8`-RC4uCeIi+Ndp( zRGcqe796gCBu*fQ)4ZMS-D9>%gXW3dl#+4OG2s5{j`>wdlg9ar4>S3O^A_EM?TtEz zNr8^d;Q7Hc_c%P&}_96%fR-Q5%R5h6>A<%>L)ID5PlYMI|@rM}_#K*}$gd zZOYASqKvIWt#AxSeSLq#>fpu8_=L1hM+~E(R(x=GCWy3LnBIP3Lwu6H2TAEU z-s0@?5sdDF&1yIAOFn81Q|#w~Req0=?UQ=XKu-~3kDaiMx2MKz`siodnQnTYN`Snp zkTxb*Bz6DDv2^LsJJP;y_&f=l9LyNDFUoj+&F-3Um4q{N!@2YlPQ+Vm50!ozPGC&G z_K9SP0RDQbVmff*lQdHZ34K$B>w#!oGIl{Ahwk)4-rF}aO2{;x4MWDABvhOY#;<-T z2Rv5Z8Gh9U{Bw83ZwsOeU;2q`xsv|JiN4@E{P=$Vf!N2~vNA)UkN`}RW<&E=Tbt@9#x6(x!#nL8UxDKQXDF{ zZ#6&fp#B_Iz`CP7ZPLE;8eS3pkFz}hkT%#o?6k&DjqOli2=Tkp;XnEV`rDxnYkKg< zsE{a&J0|2m2vy)UgzlQ#9vyOzr3V*$eh7rz{_>IkDsD^nN!c81*Ia@D88?6|aq0MN z7t>?It;zRB_x_t(3|MI-zeJxu_MGjtH?MnJtUY2VdVcOnFNgoj@NqETC($+nNvA3N zt0ucX<;tw56MHvYvFJ-cy8zKnW3~jh)Zl5s?_hoKu#Ln+ zhj&!iO|l+mXJeoCs`r!(ao&D0{8540!P!I%f_Sm+?cSsc6&ZAL$2>YC@Z}CGqiYxt zRn5R$w|)8pV~^U{_~SCplpft7bL>M>JM0~Vb{U_CB4JkQGagLI`hGp$ItmMq^jX4M zvVdbO!*6c)qsaUAM3dbG4#iZpVEscu>UIx6Vf_(r)H$IMpRU?`lu;p^duK0{e_%-z zRD?z$`bpq^l?OiT+E7>1!c=6cSS{ELQ0-8keeKqs!(a_pCJ#V)27wrX^7OH}9cJ+) zjf_#69;rs85T1s!p7D4cYT6!02FgxmWOb4@?Ed_39&?Fo&npk|>KhXjaVn}k7#b+r zEZyb{r#jMXj~qORt60sSkn~$lE$xNm8vBp$6m-n4`)`2L4!h34RBIYkS&6LzpUP_~ zc%35RCVm0cU1)GOwsy9zy8Qu70zci4bVq$f3*I7i4w zkXbU6j6j~@nb&G2UIR|YE?Y=MU>zm+_)6nK&`9Zs=)Ys(i-O8$*^#6#y+WhNiAk%5 ztL!J#>)Wi^V0!KRkR;XMUJ=;v^w|Un(bobTC^{U(#cK1>MT$9nH|*hXOebDc)@VAGFj3IP`FBzx z!w{;>2&TbMm*;ZTtf3Sj)>IWeM&PhVHWOh;9{-$k?+ll}dr&s$@U93UH2Dx5m$6?s zlVMoIiwL#Kp%<)ByK$TlCwcYAH3RqAGVc6&w`W^w)O2~FgLpI{V3@x>W2pKxQI`Ck zc{?o=8yao@Jg(b)?=YgTTB}#CV$P8b4{G-F%+Q;4;yWP?IGH=Fb}&NjWCXrwW;d*E zdvRf|+f;`3EnouqJFI+GpICR42mR)&1&K2yRCwXZ2N9|y$4=pdW8qxo=>lacF4;It zu;cC}jG1dy5MS_K7dQ#$S@Aq{f5@$FbI50LNrv*q5`qDwjZ7rp_%*%=@6y~4OPMLw zEQH2(UPoG~yP;5argA~73_iQA_+>?MXnbG{MnRD-{79S9@_E)2u7DAfK}hG)P7(5! z_mDXq-=^tgI8>+1k2u1C0(4A__d~1okU75R5fH-xKi37xSM^vJcuT3b4box9>w;Lp zlc~3V?O#TDj$iOaqXaF7&Rir=G~EFZP3WkiK{6l9JL7ax&na0<0}Glze8{v880Sv) zRDI>o4->2y+C>xc;%(w-GA6M5Tzo<(55HpsU$uDrNUey7sf zJ111HU398=u<+vYBnL*DffY%q)kAZv2S*o3d(suup{&4vI3X9~;Rub0d8BX?c_ zQf?4QB<4$NoI%NyZ;*3qqU zdF?udZ?~g`Ko5RMO;~^O66!ZLl(q5b|($pSZUf3F}N)%{X zHGlB;IY$oGsE%iYz2!-7*R0+JpQ8id2r=nKup*RV%5*X0o)eBK5pX`kKr!B1VPgz- zpLAWncIZbb9LRSS-^%1uxe9YID&+VHstw243;RG@Dm0j^^jqY#{%c~{D{KN=#Lf3NTZp|RX~ZZ7)J4ZrS3t_iQa>Fn~EeX9@< z)pp)Rgwb57RN!3OOZwezlJA*vbQ`J!--l_WqWelK(gT1@OX~UDUe^6@ZMnQwj10ZI zgo)R6mkq;QDY&U8_DCheQ4U>))2K#E`vY zNMaaEWEZk#X|t4_v1N;Fh3v{w82b##J|eP3M9RL6Er~=)_I(Y*;Jnc1d_L#;e$P4A z`R{lA-hUd`#e12V=j(pm_v1DgpmxYT-%dRPDf_wFFXc^ailNP=bG~HATOK~dvq(ML zIbroWGI7u?%fqze1~+1X#8Ygl?!qZyrv1+pDO;)K$^16$sORq9bX&TUo&AkeIjZbo zEZT6`&I7ejBeI#8H=+_Ac|@U;OIz_do&5=wnMO3yVxDYr@^hRncjVW(%_|4)kNj%( zX@*G1&9ydnXe~k(WaUM5w^@9HmqgplW;^Gz_z26 zT#0Ta?U!^1lR|tm+Wm0Co{xY|?ou1*d{*ncO@wt-w%Rubs>D%7nV zD+SceKts5?^uC-@_YCuWY&PM;{;++z-Q=wDT?xvrQ26X!TBVc)uvZfP2`ev8&Km=+ zAZhLdODo)D#XoT?eS!y$WW2QPIsb#yQue_aZw;*kjp_PbZ_D9Pqshz*Ae0R9af|7E zoN-0NYawoar*5A5j8~t%A#))MX5F~sWdhvL`W>WvwA$S%Cq4n7F!065{0M;`Cg9R@ z4+?6Zih`v{tTymtYlI$8$9G2d*cJT4tq+G_y-OEFRo!m$2OeRz4##9y*@5ggKO3E^ zALj74oYYn=W8dafj)&)y6_{Y2!W2|Q|M)Dxe7sM<72ouolN;6gU?6YgOBpa)TyZcN9`6PGvwQML1wTfe{X0|B z5u_Rf<*fhJf;4l8W_+weUT=Gs?Ep%l3#`D4d;joGN9O~7<-qR@=j}^(hJ{WuMLAM- zmC>&Fxd}!stUp$gRGO|pXL#8>w7Xl8&l)lB7X4xl(50t;4)*2&l47E2P}M9TH>M1J zq@j-Pm^=a(|NKxueR%)^$$d*p;Drr(##~&ks_ELSzxk}Sv*O;HF;2lzN>Uwf_RJ+V zs$r}Z$<~S~p5COZb?1+9Iq#lCR*mE!lY6Pk}IP=D($ZL&GrX*z?aBB*J0%b79pZ5+X5`w_C;7x zWCuU@O8>8R7i49B%MJbgYb?*CL@p~@;%_nKO8kDuP!!-!33^+DIu4U{|1e?G1m(75 zog~&IVi|JICP^r4iY9$|tg4485f2^VpQC6NM&IDy6${?EY00BHy3u zP@8!)=qnB+G~;wALY2~^*%Aq@Mch> zv|_Jl%}_+fqPW)l*xCJsc+-o8JAjv=;hTC!o(?p0a$V-!335A_b8qii^0>z4vi`_vHg6#1AUe@m;jQO=D6*lV6#{@ve z?SS~eM4!u7_mzmNV?(4l;|e+bqFW*vbWVpQu+@8r{4 zlJB?0T6+U2TestAudK+$+)uQH-Fox=RvSsVH>7Tt>#AhYc9-kAZ&I#tC%7mvYCZ^hdj+91dYurnGKsHeA zS>k)X45{Qcz@O$^Ju?}apORTE;v_lTv*^h?^7(GT7kSrKc2dyL%6OvSQKEk^Jw#iB zST|amp&{1bEkMW!vn(4|IK(ftIb^aLy!%o4>M`Ir3~KmSwhJyVd^UG}pp4-t>2PJY zc?@S0T`|onmjL(5U@D%NCk;te{FypMFVJHU>s6oR%yucwf54+2~cn&E6tDD1W4%DClFDB-pM6v zR;l%wzSo4fh_F&Wh_jzID7cP9oC~0vK@QR7sggx}f2{Q_LhW8?^C6 zG*p~OG2Ml)%rq1PO1tdAF#`^UN&h-Vvw}cz@6J9}gLG8Q4~xKHpLUG zDB|21am{5nX??up3z%CSa$M4l#I$U*=4<&>15o;rg7Iy zbnE4dU?0Ic;-;x*C#7g1+VEs%fySIL@}9`3gycBf_(d_S-!zWvQ7Wy1ml89&8wyTY z1rv2Dp#XtLzGzh~U7&txMIGD}HvoLy3O(A$M?3XG-{#B}i#0j(ft;>`ZaH@)toCH< zJfPOu_3T*p-c5gYw%Cujc20Ss8r#WeR=nQ)5O2g@V)LSrAKi1dnR+Jcb>H=1cv*30 zENLPCZtlAJbl7vwhg!`z} z*uVbMF{tr${FE@k>-lyt+`DkD!gx6#v(QxrY&Nx1 z#v4q{;N)SqCL8pI5)7ajRB;R*r*5{36i=zka{p^*2$){I+w1g{+WzfWbnMLToPS4s z5R}amvD2;H%y$eK_^YG@)QcVQT^-oOyFX)8V9Z#%yc*h3K z1EP$Ep$8fNW8X=+mnR4Nf2HLfQF#9~X?Wyn{%=WkM{ejppxqnX6k1aT^8UZr7vM6K zkDzJVMnEe!oLwM573$MSc`bw^(0=F5HX$Nr!6LLWH>A|3*rxw6&!5JWb=4(!CnXL5 z4@M{(W=-PnS^xj~+Ohc$YpX;&TG5$30+MM#g}1wq9b6v@N&ypfGvMh}58cYB_bnLn zq#4QY-2kz$ddtOwUU>4YJLia}@*|8s>v_yKQsrglEtPw$A5bigPCdOfV*sz)NcjQv zfS|$^>9>zsK`osL*;_H)x=4s^Z9WF~z3a8o#-{(DZo?)(-;JUWTjP zOz%;2C4e^*wO(@)?@n1>tKo89i^c3QgWJE!Qy-s&393c8rvb8;5wFMtXIc z`qd4+fRzT0*A(4Y(G*GO5sA`;bY*ANrsR1UgO>}I>K z5D2DAKrOB!Z?TjB)ypJ{+a_(P=ZPG6PrfhU5J{5(F>E~PPxQi{CDzleY(wE@%f`gh0RJ1A0Yz5BoohMwuk~xMCL-evq3E(N=h|3iylU7{Mr{ zSL#XU!ie?)%j8+F?&-5vm5l7rH9+IgUm>*q+wYrj;xLu3t!>(PQock|_RV{h&GSAH z%Lwl+{=84F(i5$5$>;9JOW!G14eSfNo=kLg(Q0&NC*}ev&eT5U7VovWXN^}K*ARKS zv?8iSdr>kI@(eM$DXW(6ri|_*^i%KWv2M0lb+@?J`7Q*s9~iC}ahiFAi;Wib3f zPwqXhUs=fP)eFkM+`InCgB62!D?x(p577(hW3oh$GLA4kxGWQkm3~IJZ=H1>48iZD z5fcJ)&i~bj+vkK6=I+h#`epru#0kmrAf`hg3sCO)5RJL8)R#!jJDA$x9vgYGGi)AJ zAUw{s~sWm(eHhV3tSZwZnLz3SjNfG=FT zEUL|Afo%}@OslvwSRV?1{Bq!wHkZj4lqB#G!sMKW8j@Adg_hXZ*}#w3@B?KF=0r~7UBR==K)RQVsrdPZ~2>)#rTpAfsIGxfyQvcyGwv)$;! zT4x`pSIK-Fx)!6)u07`2hm{fBpm&c%C-2L!-ULt&gPjDw;5`e&t5D-k`Zbp`%~5)KM9GUE9U3p8E^0+55W zq6UUUO=nYsg7~#!w}Du{)@iXpOQ7e;;L_R6=ow8KT`wUU%HE}rrHDx86K|PCJRmrL zqTZ`q+0F|oq%RfSr3+GA?>;3?7Qu(cBQi=0Ex|w6YxM1NeL|~O;<$CTe90*N=gP)X z*&p28)9>Ef-oR8Ar4`3RR*y4Z#$;7-aSB=dtJ7&ex_o&sW>B^==qF#4GqX{c7d$XXjH( z<6Q-givY*L<(g}Caqjr*3uuRAC~^C0u#jN9IT`crw#&Lki7G*SbfYp;KLOw#ZX!od zC=r#Rmz;F#BTg$l9YMQ)RI;Ed#`!5kAddr+N4@Y(oJW}H>L8cRrZE#W0+Xn7sQG|O_(&qh|S2XHXCaq{Z z*r!T5peuApkP&N6qAidKx-MN@djIe)#e(mgOU@tJO(@~GU0CjTjXZARV(6J-22A)60MV6~ zm>iff^=NBM=f|Y0tpvj(VK(hut0Xd|eb%{!Qj@zmDo^0W+lgKDI6Lqh5-gsw0+d+# z5lZa&&sy!-6CVHRz!1yO&ll6RJDK_w-?j?27>wq`EU=P?xARg;;;(gq93WTrT0FJE zPOLsnR~He%o<1@OI~BK~R@bzl>}4x{Qj@3Y;s-E&ICncW*SU8ay(*>6>8) z)1fdXuuhjU_Nu(w=sFV^F4MkYnj9=cTp4XBhQ$fdeBtljY3%L^1DSXZ5P^&nQ*h27$B1 z67h75a83hyti>xQW*C&y%Kxw}0RQRwFz^qztQ&}Hok_F<;~_&M|Gv8HxWM(szvvG~ z-D*Q19G2FaGQJb>xZ@Jo{6{iDbyt~Qj)8grFdt6YgV}Ps1TOEHw#@Sok$!oLn_a;> zZHL>?&oa#h`zC#VUF^xSC5JE%aU_~EGNhX!o%{bc(GvGQZGT_garnRRmvs=n0b{)2 zhlfpyM-*b~WRR*G`@4ycvpF~jUinvR1`c%GQ5F9UgMc%i)P6Vlr(Dkvn11!V{?;w8 zmV${^@F*TKG)t$+M)n+AmV9sFKt<%<+y11sgLY{uD`UA644ZRvn5Qk5X)_Oz4|l%< z0K`WheY+%2MM2>^svmx+==Hi(Pr^)S0_Ibx)k(>~3UWVw^yNPQ4*oFnVL+p`90KIW zsnf0jH7IHm_Cf+h(Ug20UIG4vXVXAf$qXfnm%ANQ-AL_)m&OoMF{8M#F4gXFq5AxM z;guaa6F_jl9+4Z?*B-F&N+@Q~ql(1BM=a z@6#1hbOW_*vSi9@rk6T=t|+U-KAe%NX(Rakfcyseue&;L1L6Lv_3_bA6#W zvpN1j0Exqnrz5%pMO==tfA3q`mO z{uOe3PYa%Rs{?ke7(tO|orV2y$DRLh$DQE>0LF5l0r`#B(R(E=vx5}Elq03+<-k*) z5FXSN2H!+FfEi5rRqi8u<%1xt$D<>&EJ|A!yP|fZ$D0Ed8W1;_#dWU3a>{8)Jp&!nq-ZD%9TR>68mN z@KMo0uAA~O|0XCX>%ID?pzMD9xyEoR-iGHG&l03*o}~ zqtfjV{fbH-U^+0GXKmERlO$CplS|%k@bE~f= zD4&cz8CQT0i^}=v|6YO)miqw#P^|e&`Wc9uu8L{A^(Is9n){pG8CSZi+dzJ-a`|!o zn)%dv^$(HpjAW^;+SkjrkhH4 zmL4L!qO_9y&dw2LT7HQ`t2bB4bU*C`4cK$yJ}Z>{;t@Nn%=(26;HW>^6Sg|u6IQC( zw~UA`N27b?%oK~iah^4j0^G7f67>`r`+QVzV&4zyT-NZ?arTVX-8kG(l}AQm)#00d z1T4KQhw>b~oW31$tB&L?Ie$-^HCLr8&qzDf_E=!Teh;s|0@$iwkWh>|^~sgSs|Xv) zv9KZn7xt~IU*p3DM*4=+j>}u=sbMugu8x=N-jcU~+pSnlK=q2j z{;b*G+kiR<3}Bez{xI~*4hxZ1{^lpeSOTHS_%FH2#;hkz5}I=VpRbRNKo3gaUW*hv z+E4|5;lBy<{2$w;cJR~b^=*F+J{geD`ExwI0)(eOhh+>Xx;;L!S*<~#!STUI{=ffv c-fuC&fnopU4MCVAH4!*8RIjTPDOm*n2S(|TlK=n! literal 0 HcmV?d00001 From f50a4e0b68ed8c52f4529828e245e82efa992eaa Mon Sep 17 00:00:00 2001 From: Akshay Shah Date: Tue, 27 Sep 2016 19:22:10 -0400 Subject: [PATCH 10/10] updated readme --- README.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/README.md b/README.md index 61260d9..0a49626 100644 --- a/README.md +++ b/README.md @@ -16,9 +16,9 @@ Max size: 230 CUDA Summary -![](images/GPUCUDAstats.png) +![](images/GPUCUDAstats.PNG) -![](images/cudaCoreStats.png) +![](images/cudaCoreStats.PNG) Memory I/O looks to be the bottleneck over the GPU, meaning changing the variables over to shared memory might make a difference. (not checked, just speculating)