From e441f6a7b20bc9f149d857baa4171ad127afa29b Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Sat, 19 Sep 2020 22:15:17 -0400 Subject: [PATCH 01/10] finished gpu naive scan --- src/main.cpp | 270 ++++++++++++++++++------------------- stream_compaction/cpu.cu | 42 +++++- stream_compaction/naive.cu | 92 +++++++++++++ 3 files changed, 265 insertions(+), 139 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 896ac2b..6ac57b8 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -15,140 +15,140 @@ const int SIZE = 1 << 8; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two -int *a = new int[SIZE]; -int *b = new int[SIZE]; -int *c = new int[SIZE]; +int* a = new int[SIZE]; +int* b = new int[SIZE]; +int* c = new int[SIZE]; int main(int argc, char* argv[]) { - // Scan tests - - printf("\n"); - printf("****************\n"); - printf("** SCAN TESTS **\n"); - printf("****************\n"); - - genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case - a[SIZE - 1] = 0; - printArray(SIZE, a, true); - - // initialize b using StreamCompaction::CPU::scan you implement - // We use b for further comparison. Make sure your StreamCompaction::CPU::scan is correct. - // At first all cases passed because b && c are all zeroes. - zeroArray(SIZE, b); - printDesc("cpu scan, power-of-two"); - StreamCompaction::CPU::scan(SIZE, b, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(SIZE, b, true); - - zeroArray(SIZE, c); - printDesc("cpu scan, non-power-of-two"); - StreamCompaction::CPU::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(NPOT, b, true); - printCmpResult(NPOT, b, c); - - zeroArray(SIZE, c); - printDesc("naive scan, power-of-two"); - StreamCompaction::Naive::scan(SIZE, c, a); - printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); - printCmpResult(SIZE, b, c); - - /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan - onesArray(SIZE, c); - printDesc("1s array for finding bugs"); - StreamCompaction::Naive::scan(SIZE, c, a); - printArray(SIZE, c, true); */ - - zeroArray(SIZE, c); - printDesc("naive scan, non-power-of-two"); - StreamCompaction::Naive::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); - printCmpResult(NPOT, b, c); - - zeroArray(SIZE, c); - printDesc("work-efficient scan, power-of-two"); - StreamCompaction::Efficient::scan(SIZE, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); - printCmpResult(SIZE, b, c); - - zeroArray(SIZE, c); - printDesc("work-efficient scan, non-power-of-two"); - StreamCompaction::Efficient::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(NPOT, c, true); - printCmpResult(NPOT, b, c); - - zeroArray(SIZE, c); - printDesc("thrust scan, power-of-two"); - StreamCompaction::Thrust::scan(SIZE, c, a); - printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); - printCmpResult(SIZE, b, c); - - zeroArray(SIZE, c); - printDesc("thrust scan, non-power-of-two"); - StreamCompaction::Thrust::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(NPOT, c, true); - printCmpResult(NPOT, b, c); - - printf("\n"); - printf("*****************************\n"); - printf("** STREAM COMPACTION TESTS **\n"); - printf("*****************************\n"); - - // Compaction tests - - genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case - a[SIZE - 1] = 0; - printArray(SIZE, a, true); - - int count, expectedCount, expectedNPOT; - - // initialize b using StreamCompaction::CPU::compactWithoutScan you implement - // We use b for further comparison. Make sure your StreamCompaction::CPU::compactWithoutScan is correct. - zeroArray(SIZE, b); - printDesc("cpu compact without scan, power-of-two"); - count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - expectedCount = count; - printArray(count, b, true); - printCmpLenResult(count, expectedCount, b, b); - - zeroArray(SIZE, c); - printDesc("cpu compact without scan, non-power-of-two"); - count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - expectedNPOT = count; - printArray(count, c, true); - printCmpLenResult(count, expectedNPOT, b, c); - - zeroArray(SIZE, c); - printDesc("cpu compact with scan"); - count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(count, c, true); - printCmpLenResult(count, expectedCount, b, c); - - zeroArray(SIZE, c); - printDesc("work-efficient compact, power-of-two"); - count = StreamCompaction::Efficient::compact(SIZE, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(count, c, true); - printCmpLenResult(count, expectedCount, b, c); - - zeroArray(SIZE, c); - printDesc("work-efficient compact, non-power-of-two"); - count = StreamCompaction::Efficient::compact(NPOT, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(count, c, true); - printCmpLenResult(count, expectedNPOT, b, c); - - system("pause"); // stop Win32 console from closing on exit - delete[] a; - delete[] b; - delete[] c; -} + // Scan tests + + printf("\n"); + printf("****************\n"); + printf("** SCAN TESTS **\n"); + printf("****************\n"); + + genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case + a[SIZE - 1] = 0; + printArray(SIZE, a, true); + + // initialize b using StreamCompaction::CPU::scan you implement + // We use b for further comparison. Make sure your StreamCompaction::CPU::scan is correct. + // At first all cases passed because b && c are all zeroes. + zeroArray(SIZE, b); + printDesc("cpu scan, power-of-two"); + StreamCompaction::CPU::scan(SIZE, b, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + printArray(SIZE, b, true); + + zeroArray(SIZE, c); + printDesc("cpu scan, non-power-of-two"); + StreamCompaction::CPU::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + printArray(NPOT, b, true); + printCmpResult(NPOT, b, c); + + zeroArray(SIZE, c); + printDesc("naive scan, power-of-two"); + StreamCompaction::Naive::scan(SIZE, c, a); + printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + printCmpResult(SIZE, b, c); + + /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan + onesArray(SIZE, c); + printDesc("1s array for finding bugs"); + StreamCompaction::Naive::scan(SIZE, c, a); + printArray(SIZE, c, true); */ + + zeroArray(SIZE, c); + printDesc("naive scan, non-power-of-two"); + StreamCompaction::Naive::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + printCmpResult(NPOT, b, c); + + zeroArray(SIZE, c); + printDesc("work-efficient scan, power-of-two"); + StreamCompaction::Efficient::scan(SIZE, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + printCmpResult(SIZE, b, c); + + zeroArray(SIZE, c); + printDesc("work-efficient scan, non-power-of-two"); + StreamCompaction::Efficient::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(NPOT, c, true); + printCmpResult(NPOT, b, c); + + zeroArray(SIZE, c); + printDesc("thrust scan, power-of-two"); + StreamCompaction::Thrust::scan(SIZE, c, a); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(SIZE, c, true); + printCmpResult(SIZE, b, c); + + zeroArray(SIZE, c); + printDesc("thrust scan, non-power-of-two"); + StreamCompaction::Thrust::scan(NPOT, c, a); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(NPOT, c, true); + printCmpResult(NPOT, b, c); + + printf("\n"); + printf("*****************************\n"); + printf("** STREAM COMPACTION TESTS **\n"); + printf("*****************************\n"); + + // Compaction tests + + genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case + a[SIZE - 1] = 0; + printArray(SIZE, a, true); + + int count, expectedCount, expectedNPOT; + + // initialize b using StreamCompaction::CPU::compactWithoutScan you implement + // We use b for further comparison. Make sure your StreamCompaction::CPU::compactWithoutScan is correct. + zeroArray(SIZE, b); + printDesc("cpu compact without scan, power-of-two"); + count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + expectedCount = count; + printArray(count, b, true); + printCmpLenResult(count, expectedCount, b, b); + + zeroArray(SIZE, c); + printDesc("cpu compact without scan, non-power-of-two"); + count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + expectedNPOT = count; + printArray(count, c, true); + printCmpLenResult(count, expectedNPOT, b, c); + + zeroArray(SIZE, c); + printDesc("cpu compact with scan"); + count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + printArray(count, c, true); + printCmpLenResult(count, expectedCount, b, c); + + zeroArray(SIZE, c); + printDesc("work-efficient compact, power-of-two"); + count = StreamCompaction::Efficient::compact(SIZE, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(count, c, true); + printCmpLenResult(count, expectedCount, b, c); + + zeroArray(SIZE, c); + printDesc("work-efficient compact, non-power-of-two"); + count = StreamCompaction::Efficient::compact(NPOT, c, a); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + //printArray(count, c, true); + printCmpLenResult(count, expectedNPOT, b, c); + + system("pause"); // stop Win32 console from closing on exit + delete[] a; + delete[] b; + delete[] c; +} \ No newline at end of file diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..0d06ff8 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -18,9 +18,17 @@ namespace StreamCompaction { * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); + // Compute exclusive prefix sum + // timer().startCpuTimer(); // TODO - timer().endCpuTimer(); + if (n < 0) { + return; + } + odata[0] = 0; + for (int i = 1; i < n; i++) { + odata[i] = odata[i - 1] + idata[i - 1]; + } + // timer().endCpuTimer(); } /** @@ -31,8 +39,15 @@ namespace StreamCompaction { int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int o = 0; + for (int i = 0; i < n; i++) { + if (idata[i] != 0) { + odata[o] = idata[i]; + o++; + } + } timer().endCpuTimer(); - return -1; + return o; } /** @@ -43,8 +58,27 @@ namespace StreamCompaction { int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int* bdata = new int[n]; + for (int i = 0; i < n; i++) { + bdata[i] = idata[i] != 0 ? 1 : 0; + } + + int* sdata = new int[n]; + scan(n, sdata, bdata); + + int o = 0; + for (int i = 0; i < n; i++) { + if (bdata[i] == 1) { + odata[o] = idata[i]; + o++; + } + } + + delete[] bdata; + delete[] sdata; + timer().endCpuTimer(); - return -1; + return o; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..6289d66 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -2,6 +2,8 @@ #include #include "common.h" #include "naive.h" +#include +#include namespace StreamCompaction { namespace Naive { @@ -16,9 +18,99 @@ namespace StreamCompaction { /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ + + + /* Copy initial data over and pad 0's if out of scope of initial size + * aka the input array has a smaller initial size than the final array, + * and anything larger than index [size of input array] will be 0 in the output array + */ + __global__ void formatInitData(int initSize, int finalSize, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= initSize && index < finalSize) { + data[index] = 0; + } + } + + __global__ void add(int n, int ignoreIndexCount, int* odata, const int* idata) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index < ignoreIndexCount) { + odata[index] = idata[index]; + } else if (index < n) { + int x1 = idata[index - ignoreIndexCount]; + int x2 = idata[index]; + odata[index] = x1 + x2; + } + } + + __global__ void shiftRight(int n, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + if (index == 0) { + data[index] = 0; + } else { + data[index] = data[index - 1]; + } + } + + // Careful with non-power of 2 void scan(int n, int *odata, const int *idata) { timer().startGpuTimer(); // TODO + if (n < 1) { + return; + } + // Calculate the number of elements the input can be treated as an array with a power of two elements + int kernelInvokeCount = ilog2ceil(n); + int n2 = pow(2, kernelInvokeCount); + + int blockSize = 128; + dim3 blockCount((n2 + blockSize - 1) / blockSize); + + // Declare data to be on the gpu + int* dev_odata; + int* dev_tdata; + std::unique_ptr tdata{ new int[n2] }; + + // Allocate data to be on the gpu + cudaMalloc((void**)&dev_odata, n2 * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed!"); + + cudaMalloc((void**)&dev_tdata, n2 * sizeof(int)); + checkCUDAError("cudaMalloc dev_tdata failed!"); + + // Transfer data from cpu to gpu + cudaMemcpy(dev_odata, idata, sizeof(int) * n2, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy dev_odata failed!"); + + // Format input data (pad 0s to the closest power of two elements, inclusively) + formatInitData << > > (n, n2, dev_odata); + + std::cout << "kernel invoke count: " << kernelInvokeCount << std::endl; + + for (int i = 1; i <= kernelInvokeCount; i++) { + int ignoreIndexCount = pow(2, i - 1); + add << > > (n2, ignoreIndexCount, dev_tdata, dev_odata); + + int* temp = dev_tdata; + dev_tdata = dev_odata; + dev_odata = temp; + } + + // Shift things to the right to make the inclusive can into exclusive scan + shiftRight << > > (n, dev_odata); + + // Transfer data from gpu to cpu + cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_odata failed!"); + + cudaFree(dev_odata); + checkCUDAError("cudaFree dev_odata failed!"); + cudaFree(dev_tdata); + checkCUDAError("cudaFree dev_tdata failed!"); + + // Calculate the number of blocks and threads per block timer().endGpuTimer(); } } From bb2d77386036a7979dc8d4351c7034cf2792947d Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Sun, 20 Sep 2020 17:30:08 -0400 Subject: [PATCH 02/10] finished efficient stream compaction --- src/main.cpp | 14 +++ stream_compaction/common.cu | 16 +++- stream_compaction/efficient.cu | 152 ++++++++++++++++++++++++++++++++- 3 files changed, 177 insertions(+), 5 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 6ac57b8..645aa2a 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -29,6 +29,13 @@ int main(int argc, char* argv[]) { genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; + /* a[0] = 1; + a[1] = 4; + a[2] = 9; + a[3] = 0; + a[4] = 2; + a[5] = 1; + a[6] = 4;*/ printArray(SIZE, a, true); // initialize b using StreamCompaction::CPU::scan you implement @@ -104,6 +111,13 @@ int main(int argc, char* argv[]) { genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; + /* a[0] = 1; + a[1] = 4; + a[2] = 9; + a[3] = 0; + a[4] = 2; + a[5] = 1; + a[6] = 4;*/ printArray(SIZE, a, true); int count, expectedCount, expectedNPOT; diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..98e98e1 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -23,7 +23,12 @@ namespace StreamCompaction { * which map to 0 will be removed, and elements which map to 1 will be kept. */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { - // TODO + // TODO + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + bools[index] = idata[index] != 0? 1 : 0; } /** @@ -32,7 +37,14 @@ namespace StreamCompaction { */ __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { - // TODO + // TODO + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + if (bools[index] == 1) { + odata[indices[index]] = idata[index]; + } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..6d43311 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -2,6 +2,7 @@ #include #include "common.h" #include "efficient.h" +#include namespace StreamCompaction { namespace Efficient { @@ -12,13 +13,102 @@ namespace StreamCompaction { return timer; } + /* Copy initial data over and pad 0's if out of scope of initial size + * aka the input array has a smaller initial size than the final array, + * and anything larger than index [size of input array] will be 0 in the output array + */ + __global__ void formatInitData(int initSize, int finalSize, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= initSize && index < finalSize) { + data[index] = 0; + } + } + + __global__ void upSweep(int n, int offset, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index < n && index % offset == 0) { + data[index + offset - 1] = data[index + offset / 2 - 1] + data[index + offset - 1]; + } + } + + __global__ void downSweep(int n, int offset, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index < n && index % offset == 0) { + int halfOffset = offset / 2; // Helps to find left child + int t = data[index + halfOffset - 1]; + // Set right child to be the same as parent's value + data[index + halfOffset - 1] = data[index + offset - 1]; + // Set left child to be the sum of parent and parent's sibling + data[index + offset - 1] += t; + } + } + + __global__ void setLastElementZero(int n, int* data) { + data[n - 1] = 0; + } + + __global__ void shiftRight(int n, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + if (index == 0) { + data[index] = 0; + } + else { + data[index] = data[index - 1]; + } + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); + //timer().startGpuTimer(); // TODO - timer().endGpuTimer(); + if (n < 1) { + return; + } + // Calculate the number of elements the input can be treated as an array with a power of two elements + int kernelInvokeCount = ilog2ceil(n); + int n2 = pow(2, kernelInvokeCount); + + int blockSize = 128; + dim3 blockCount((n2 + blockSize - 1) / blockSize); + + // Declare, allocate, and transfer data on gpu from cpu + int* dev_odata; + cudaMalloc((void**)&dev_odata, n2 * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed!"); + cudaMemcpy(dev_odata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy dev_odata failed!"); + + // Format input data (pad 0s to the closest power of two elements, inclusively) + formatInitData << > > (n, n2, dev_odata); + + for (int i = 0; i <= kernelInvokeCount; i++) { + int offset = pow(2, i + 1); + upSweep << > > (n2, offset, dev_odata); + } + + setLastElementZero << > > (n2, dev_odata); + + for (int i = kernelInvokeCount - 1; i >= 0; i--) { + int offset = pow(2, i + 1); + downSweep << > > (n2, offset, dev_odata); + // Transfer data from gpu to cpu + cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_odata failed!"); + } + + // Transfer data from gpu to cpu + cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_odata failed!"); + + cudaFree(dev_odata); + checkCUDAError("cudaFree dev_odata failed!"); + + //timer().endGpuTimer(); } /** @@ -32,9 +122,65 @@ namespace StreamCompaction { */ int compact(int n, int *odata, const int *idata) { timer().startGpuTimer(); + // TODO + + int blockSize = 128; + dim3 blockCount((n + blockSize - 1) / blockSize); + + // Declare, allocate memory in GPU and transfer memory from CPU to GPU + int* dev_idata; + int* dev_bools; + int* dev_indices; + int* dev_odata; + + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_idata failed!"); + + cudaMalloc((void**)&dev_bools, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_bools failed!"); + + cudaMalloc((void**)&dev_indices, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_indices failed!"); + + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed!"); + + cudaMemcpy(dev_idata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy dev_idata failed!"); + + // Compute temorary arrray containing + // 1 if corresponding element meets criteria (of not equal to 0) + // 0 if element does not meete criteria (of not equal to 0) + StreamCompaction::Common::kernMapToBoolean << > > (n, dev_bools, dev_idata); + + // Run exclusive scan on temporary array + scan(n, dev_indices, dev_bools); + StreamCompaction::Common::kernScatter << > > (n, dev_odata, dev_idata, dev_bools, dev_indices); + + std::unique_ptr bools{ new int[n] }; + std::unique_ptr indices{ new int[n] }; + + cudaMemcpy(bools.get(), dev_bools, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_bools failed!"); + + cudaMemcpy(indices.get(), dev_indices, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_bools failed!"); + + cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_odata failed!"); + + // Clean up + cudaFree(dev_idata); + cudaFree(dev_bools); + cudaFree(dev_indices); + cudaFree(dev_odata); + + int remaining = bools[n - 1] == 1? indices[n - 1] : indices[n - 1] - 1; + remaining++; + timer().endGpuTimer(); - return -1; + return remaining; } } } From 564ec5a8f85c148609e2d4ea45ccf49ea8963beb Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Tue, 22 Sep 2020 14:43:44 -0400 Subject: [PATCH 03/10] done --- stream_compaction/thrust.cu | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..df6d4e6 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -2,7 +2,7 @@ #include #include #include -#include +#include #include "common.h" #include "thrust.h" @@ -22,6 +22,14 @@ namespace StreamCompaction { // 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::host_vector h_in(idata, idata + n); + thrust::device_vector dv_in = h_in; + thrust::device_vector dv_out(n); + thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + + odata = thrust::copy(dv_out.begin(), dv_out.end(), odata); + timer().endGpuTimer(); } } From 4ca356a59aa28c8e53b34b2524d06b369b91167c Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Tue, 22 Sep 2020 21:57:43 -0400 Subject: [PATCH 04/10] finished radix sort --- src/main.cpp | 19 +++- stream_compaction/CMakeLists.txt | 2 + stream_compaction/common.cu | 23 +++++ stream_compaction/common.h | 4 + stream_compaction/efficient.cu | 42 ++++----- stream_compaction/naive.cu | 32 +++---- stream_compaction/radix.cu | 143 +++++++++++++++++++++++++++++++ stream_compaction/radix.h | 11 +++ 8 files changed, 229 insertions(+), 47 deletions(-) create mode 100644 stream_compaction/radix.cu create mode 100644 stream_compaction/radix.h diff --git a/src/main.cpp b/src/main.cpp index 645aa2a..250fd55 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -11,9 +11,10 @@ #include #include #include +#include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1 << 8; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int* a = new int[SIZE]; int* b = new int[SIZE]; @@ -161,6 +162,22 @@ int main(int argc, char* argv[]) { //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); + printf("\n"); + printf("*****************************\n"); + printf("** RADIX SORT TESTS **\n"); + printf("*****************************\n"); + + genArray(SIZE - 1, a, 50); + printArray(SIZE, a, true); + std::copy(a, a + SIZE, b); + std::sort(b, b + SIZE); + StreamCompaction::Radix::sort(SIZE, c, a); + std::cout << "Correctly sorted:" << std::endl; + printArray(SIZE, b, true); + std::cout << "Yours:" << std::endl; + printArray(SIZE, c, true); + printCmpResult(SIZE, b, c); + system("pause"); // stop Win32 console from closing on exit delete[] a; delete[] b; diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index 567795b..eb71175 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -4,6 +4,7 @@ set(headers "naive.h" "efficient.h" "thrust.h" + "radix.h" ) set(sources @@ -12,6 +13,7 @@ set(sources "naive.cu" "efficient.cu" "thrust.cu" + "radix.cu" ) list(SORT headers) diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 98e98e1..8c99a0e 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -47,5 +47,28 @@ namespace StreamCompaction { } } + __global__ void shiftRight(int n, int* idata, int* odata) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + if (index == 0) { + odata[index] = 0; + } + else { + odata[index] = idata[index - 1]; + } + } + + /* Copy initial data over and pad 0's if out of scope of initial size + * aka the input array has a smaller initial size than the final array, + * and anything larger than index [size of input array] will be 0 in the output array + */ + __global__ void formatInitData(int initSize, int finalSize, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= initSize && index < finalSize) { + data[index] = 0; + } + } } } diff --git a/stream_compaction/common.h b/stream_compaction/common.h index d2c1fed..acf6a29 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -37,6 +37,10 @@ namespace StreamCompaction { __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices); + __global__ void shiftRight(int n, int* idata, int* odata); + + __global__ void formatInitData(int initSize, int finalSize, int* data); + /** * This class is used for timing the performance * Uncopyable and unmovable diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 6d43311..98fb605 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -13,15 +13,21 @@ namespace StreamCompaction { return timer; } - /* Copy initial data over and pad 0's if out of scope of initial size - * aka the input array has a smaller initial size than the final array, - * and anything larger than index [size of input array] will be 0 in the output array - */ - __global__ void formatInitData(int initSize, int finalSize, int* data) { - int index = (blockIdx.x * blockDim.x) + threadIdx.x; - if (index >= initSize && index < finalSize) { - data[index] = 0; + __global__ void upSweepEfficient(int n, int offset, int sDataSize, int* data) { + extern __shared__ int sData[]; + int tIdx = threadIdx.x; + sData[tIdx] = data[blockIdx.x * blockDim.x + tIdx]; + for (int stride = blockDim.x / 2; stride > 0; stride /= 2) { + if (tIdx < stride) { + __syncthreads(); + int s1 = sData[2 * tIdx]; + int s2 = sData[2 * tIdx + 1]; + __syncthreads(); + sData[tIdx] = s1 + s2; + } } + __syncthreads(); + data[blockIdx.x * blockDim.x + tIdx] = sData[sDataSize - tIdx]; } __global__ void upSweep(int n, int offset, int* data) { @@ -47,18 +53,6 @@ namespace StreamCompaction { data[n - 1] = 0; } - __global__ void shiftRight(int n, int* data) { - int index = (blockIdx.x * blockDim.x) + threadIdx.x; - if (index >= n) { - return; - } - if (index == 0) { - data[index] = 0; - } - else { - data[index] = data[index - 1]; - } - } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. @@ -71,7 +65,7 @@ namespace StreamCompaction { } // Calculate the number of elements the input can be treated as an array with a power of two elements int kernelInvokeCount = ilog2ceil(n); - int n2 = pow(2, kernelInvokeCount); + int n2 = (int) pow(2, kernelInvokeCount); int blockSize = 128; dim3 blockCount((n2 + blockSize - 1) / blockSize); @@ -84,17 +78,17 @@ namespace StreamCompaction { checkCUDAError("cudaMemcpy dev_odata failed!"); // Format input data (pad 0s to the closest power of two elements, inclusively) - formatInitData << > > (n, n2, dev_odata); + StreamCompaction::Common::formatInitData << > > (n, n2, dev_odata); for (int i = 0; i <= kernelInvokeCount; i++) { - int offset = pow(2, i + 1); + int offset = (int) pow(2, i + 1); upSweep << > > (n2, offset, dev_odata); } setLastElementZero << > > (n2, dev_odata); for (int i = kernelInvokeCount - 1; i >= 0; i--) { - int offset = pow(2, i + 1); + int offset = (int) pow(2, i + 1); downSweep << > > (n2, offset, dev_odata); // Transfer data from gpu to cpu cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 6289d66..c917626 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -42,18 +42,6 @@ namespace StreamCompaction { } } - __global__ void shiftRight(int n, int* data) { - int index = (blockIdx.x * blockDim.x) + threadIdx.x; - if (index >= n) { - return; - } - if (index == 0) { - data[index] = 0; - } else { - data[index] = data[index - 1]; - } - } - // Careful with non-power of 2 void scan(int n, int *odata, const int *idata) { timer().startGpuTimer(); @@ -70,36 +58,36 @@ namespace StreamCompaction { // Declare data to be on the gpu int* dev_odata; - int* dev_tdata; + int* dev_idata; std::unique_ptr tdata{ new int[n2] }; // Allocate data to be on the gpu cudaMalloc((void**)&dev_odata, n2 * sizeof(int)); checkCUDAError("cudaMalloc dev_odata failed!"); - cudaMalloc((void**)&dev_tdata, n2 * sizeof(int)); + cudaMalloc((void**)&dev_idata, n2 * sizeof(int)); checkCUDAError("cudaMalloc dev_tdata failed!"); // Transfer data from cpu to gpu - cudaMemcpy(dev_odata, idata, sizeof(int) * n2, cudaMemcpyHostToDevice); - checkCUDAError("cudaMemcpy dev_odata failed!"); + cudaMemcpy(dev_idata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy dev_idata failed!"); // Format input data (pad 0s to the closest power of two elements, inclusively) - formatInitData << > > (n, n2, dev_odata); + StreamCompaction::Common::formatInitData << > > (n, n2, dev_idata); std::cout << "kernel invoke count: " << kernelInvokeCount << std::endl; for (int i = 1; i <= kernelInvokeCount; i++) { int ignoreIndexCount = pow(2, i - 1); - add << > > (n2, ignoreIndexCount, dev_tdata, dev_odata); + add << > > (n2, ignoreIndexCount, dev_odata, dev_idata); - int* temp = dev_tdata; - dev_tdata = dev_odata; + int* temp = dev_idata; + dev_idata = dev_odata; dev_odata = temp; } // Shift things to the right to make the inclusive can into exclusive scan - shiftRight << > > (n, dev_odata); + StreamCompaction::Common::shiftRight<< > > (n, dev_idata, dev_odata); // Transfer data from gpu to cpu cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); @@ -107,7 +95,7 @@ namespace StreamCompaction { cudaFree(dev_odata); checkCUDAError("cudaFree dev_odata failed!"); - cudaFree(dev_tdata); + cudaFree(dev_idata); checkCUDAError("cudaFree dev_tdata failed!"); // Calculate the number of blocks and threads per block diff --git a/stream_compaction/radix.cu b/stream_compaction/radix.cu new file mode 100644 index 0000000..4409c0e --- /dev/null +++ b/stream_compaction/radix.cu @@ -0,0 +1,143 @@ +#include +#include +#include "common.h" +#include "efficient.h" +#include + +namespace StreamCompaction { + namespace Radix { + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; + } + + __global__ void computeBoolsForBits(int n, int bit, int* data, int* bools0, int* bools1) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + int bitVal = (data[index] >> bit) & 1; + bools0[index] = 1 - bitVal; + bools1[index] = bitVal; + } + + __global__ void computeTotalFalses(int n, int* totalFalses, int* bools0, int* falseAddreses) { + totalFalses[0] = bools0[n - 1] + falseAddreses[n - 1]; + } + + __global__ void computeAddressForTrueKeys(int n, const int* totalFalses, int* trueAddresses, const int* falseAddreses) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + trueAddresses[index] = index - falseAddreses[index] + totalFalses[0]; + } + + __global__ void computeAddresses(int n, int* bools1, const int* trueAddresses, const int* falseAddresses, int* allAddresses) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + allAddresses[index] = bools1[index] > 0 ? trueAddresses[index] : falseAddresses[index]; + } + + __global__ void scatter(int n, int* odata, int* idata, int* addresses) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + odata[addresses[index]] = idata[index]; + } + + void sort(int n, int* odata, const int* idata) { + // TODO + if (n < 1) { + return; + } + int blockSize = 128; + dim3 blockCount((n + blockSize - 1) / blockSize); + + int* dev_idata; + int* dev_odata; + int* dev_bools0; + int* dev_bools1; + int* dev_falseAddresses; + int* dev_trueAddresses; + int* dev_addresses; + int* dev_totalFalses; + + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_idata failed!"); + + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed!"); + + cudaMalloc((void**)&dev_bools0, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_bools0 failed!"); + + cudaMalloc((void**)&dev_bools1, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_bools1 failed!"); + + cudaMalloc((void**)&dev_falseAddresses, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_falseAddresses failed!"); + + cudaMalloc((void**)&dev_trueAddresses, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_trueAddresses failed!"); + + cudaMalloc((void**)&dev_addresses, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_addresses failed!"); + + cudaMalloc((void**)&dev_totalFalses, 1 * sizeof(int)); + checkCUDAError("cudaMalloc dev_totalFalses failed!"); + + cudaMemcpy(dev_idata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy dev_idata failed!"); + + int largestNumBit = 0; + for (int i = 0; i < n; i++) { + int numBit = (int)log2(idata[i]) + 1; + if (largestNumBit < numBit) { + largestNumBit = numBit; + } + } + std::cout << "largest number of bits: " << largestNumBit << std::endl; + // Compute array which is true/false for bit n + for (int i = 0; i < largestNumBit; i++) { + std::cout << "Bit: " << i << std::endl; + computeBoolsForBits << > > (n, i, dev_idata, dev_bools0, dev_bools1); + StreamCompaction::Efficient::scan(n, dev_falseAddresses, dev_bools0); + computeTotalFalses << > > (n, dev_totalFalses, dev_bools0, dev_falseAddresses); + computeAddressForTrueKeys << > > (n, dev_totalFalses, dev_trueAddresses, dev_falseAddresses); + computeAddresses << > > (n, dev_bools1, dev_trueAddresses, dev_falseAddresses, dev_addresses); + + /*cudaMemcpy(odata, dev_addresses, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_odata failed!"); + + std::cout << "Add: " << std::endl; + for (int j = 0; j < n; j++) std::cout << i << ": " << odata[j] << std::endl;*/ + + scatter << > > (n, dev_odata, dev_idata, dev_addresses); + + int* temp = dev_idata; + dev_idata = dev_odata; + dev_odata = temp; + } + + // Transfer data from gpu to cpu + cudaMemcpy(odata, dev_idata, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_odata failed!"); + + // Cleanup + cudaFree(dev_addresses); + cudaFree(dev_bools0); + cudaFree(dev_bools1); + cudaFree(dev_falseAddresses); + cudaFree(dev_idata); + cudaFree(dev_odata); + cudaFree(dev_trueAddresses); + cudaFree(dev_totalFalses); + } + } +} diff --git a/stream_compaction/radix.h b/stream_compaction/radix.h new file mode 100644 index 0000000..aa489d9 --- /dev/null +++ b/stream_compaction/radix.h @@ -0,0 +1,11 @@ +#pragma once + +#include "common.h" + +namespace StreamCompaction { + namespace Radix { + StreamCompaction::Common::PerformanceTimer& timer(); + + void sort(int n, int *odata, const int *idata); + } +} From 088c14b99b25e5db0ccc699e341b3485cdffa3f6 Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Tue, 22 Sep 2020 23:35:05 -0400 Subject: [PATCH 05/10] added basic images --- README.md | 23 ++++++++---- img/Scan runtime.png | Bin 0 -> 22268 bytes img/Stream compaction runtime.png | Bin 0 -> 19574 bytes img/main.png | Bin 0 -> 30868 bytes src/main.cpp | 2 +- stream_compaction/cpu.cu | 4 +-- stream_compaction/efficient.cu | 56 +++++++++++++++++++----------- stream_compaction/naive.cu | 12 +++---- stream_compaction/radix.cu | 16 ++++----- stream_compaction/thrust.cu | 9 +++-- 10 files changed, 74 insertions(+), 48 deletions(-) create mode 100644 img/Scan runtime.png create mode 100644 img/Stream compaction runtime.png create mode 100644 img/main.png diff --git a/README.md b/README.md index 0e38ddb..7ab4d51 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,23 @@ CUDA Stream Compaction **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2** -* (TODO) YOUR NAME HERE - * (TODO) [LinkedIn](), [personal website](), [twitter](), etc. -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) +* Thy (Tea) Tran + * [LinkedIn](https://www.linkedin.com/in/thy-tran-97a30b148/), [personal website](https://tatran5.github.io/), [email](thytran316@outlook.com) +* Tested on: Windows 10, i7-8750H @ 2.20GHz 22GB, GTX 1070 -### (TODO: Your README) +# Scan, Stream Compaction and Radix Sort +![].(images/main.png) -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +## Performance Analysis +The runtime of the naive method on GPU is worst, followed by that of CPU, then GPU work-efficient method, and finally by thrust. + +Both the naive and work-efficient method on the GPU only rely on global memory instead of ultilizing shared memory, which potentially lead to much slower runtime than thrust method (and CPU for the naive method.) Naive and work-efficient methods also do not use up the potential of warp partitioning or memory coalescing, which are features that thrust may have to reduce its runtime. + +![](images/Scan runtime.png) + +The same explanation above is applicable to the below as well. +![](images/Stream compaction runtime.png) + + +Commparisons of GPU Scan implementations: \ No newline at end of file diff --git a/img/Scan runtime.png b/img/Scan runtime.png new file mode 100644 index 0000000000000000000000000000000000000000..143eb9d9c4cddd1ca193d7dbae932f387403d94a GIT binary patch literal 22268 zcma&NbzD?y+cqo+BA}=sAX3uO($We7!Vofak8~s5rJ&M1G^j9icY_KdEiJ9W&>`LM zUJLiWpS|z*``+KVQOBwgyDGg z3i#v~+O^`+CH2p;PadnfjjyHRy2b8Rp6^btoLdf$s5cf8R~8n2w_H)vL75lAN%LM7 z25aBKBIWpCs9bXOY2b&iv?2HEqrMVe4Pf_=3@gFH#EpEU$#C;34MudNclz9kpm@rX zylGFGtXTSjJN4HFo|we_-%ZV&3n}}a9E4#1OXjK^zivYRxK{#QZ-S#;x^$Hdx;{c+ zfa^a$_iMrYug@>}V}norH=kd^#0S6p=kG6JkwBmP^*K8D70!*$lf!M3p&Vsyr^RQ7 ziy0Bdy(xmxk&*X?hleTLZ=?vjDsppkQ&1frtrVMljiiBVYJSB4M>3DVw3@1OmcCXq zVcM6*uBN6&;Y~^*>>~gAIys}T%j(6G+$-UDmq<#C`T5O__=5wtj=nz9kTzN=DXB8M z88I?(pN27Na`LWz*N;Vd_1&Af?2tU(QQ{XQB;299`Rms&-F0!#y@D}u?<1#hv5Q0d zqw)Lu`=#%#PS%zy7PGRm$9?Aw!@qVdHi}LrX@7XZ(l>z<3kxfu&;Dq~BJsiVmmQt#T4k-P z$-=HhhiR3t)FaUy`_ohBQ??LNP8XK!G(pENtND7gL;DmwrlE(YRwD&mqP9m(A3uJC zYiU`em0Jul9v;j_pg-Rb=RLI>*)Pj3D3Jg3=~K7;=xpvrtzjD;GY7}~DtJzFc1%_t z9+hd|b8j+HkBZglMjaLwmWC#2(rqR2(;<`Qku7wWO`q_KAATpZ)aNfICY; z<$1y0OplMsN0f;kv^O>;Y^=^XU*W2Sy77301taXk*j%RBvFnUwmd?q{%uMJG2@af`)SXy)_3|aflQ6NY57L2YVg?(%$6H;TJ+n!u zO)12ZqxkVE&$cITvT0Awuo~jQh$vx)Agd~dG>!=bd!8;?(bCe2oN@G2EN-io43Y3v2Fkg(RJrS7M;BvS^`(g_ z9xv0Ccpb7i-_ktVuCke+NUL^Qf^XMMcUVMFIeLpQGk2U6{8DXO&XkR@E-jIP1+zqZ zcRT_oc*KLt=DDLN(KBg&9&M`O^nTQbRf>nP36F|3DN4D&{Q2iAF_i;ix#JtpiSG#t z3KA5nJ-w#m3eUYc_cy>+lN3z2PgP#R3u8D3tqcO_>O)U0DRBkki|8EU-2MHZu778vcMj zvy%DJGt!NrA=PA?SJl^$!MChxm1hi&1BUP{tGKhe^gFbfa%!IajyB$M%Kp6 zO@9=+FEa9`3ioi&^3L>aOj!S@Rh4{1HS(HTX5Lg_HX$bvUPWXJ*!r);@poYVp9Jym zfDe8CSEl)ABKT(p`up>wS7l|~9bH|skH9VWO1v`i^5lNa&YJXpkm7b-*VQUP_*;)x z$jZTB$#eL(NlCj_$I4M9CEtE-6egm(;=6u+e$z9<%EHoZW=rsKIHkFZbmtD%l`EYM z?%Q)ixhm{?f_tlDbgUXhbCzI(LJfq?_mYon>b14ChNYVH;^GD)MS83>G&FD(l~*l7 zeFG?%4Q9S*x#cj+Se1=V%t<|zU>+qCY8rU&#!(15w4$yGc^_3UM=Z@mP>EWW=^FSP zqFP&9zt%b~81F6*&zqg97#N{jgNWu!54RVRt>18Tv;1hlEr^aL3Nd7_%q5U3y#qROiM0KB+_-O;ngTfiJo+LigX4|B(@;;W~ zh3{@h;gz|mAIF=0o7wamG%+zTjoX5W9ag_{#TS?R6oE9GH1odQaP~`>egrDJFKedH zgW~1Om&?S(M<@Di9wLS3$SUixmq3+%BkU!-DyWd*L>@B{j*+^SC5eQX7$a7Oyi0@d z4^1+wWnWb6Z%zZhEx@A@k(PQxg-^ zZDQhDV$Hq2P4FADp_$x1ohs|Zwf&J^L@l9a6BW7LOn*!Z+&}LVi?Hy7@!jUbd2+{; zuCA`BFV-E=4D-%Mo)z7>v|WRxJ9}e032RgJquCq3mR*My_Lp7N?6cdSHLrkN8+KC? zQb#_CrZ(YFn;fZF&%>o(u0L8Zbgq1?)6CA!7Jkd!L#sJfX%!#v`gPqYCCJ9oEpCLG ziHV7(KUI}Y(ZGDT&3Y-yziN}39ryiIM2v~IYxtw^t51DS!<+Ef3?jm+Hx^|POA9Lg zXqY&1*Aaz6jZq1^Z)M){y{n>m z<@qpnqo+la)XFMIssEZJ93H)9c3LlH%6UyS(_-XtTOh|f^^d<45f+DC+As4Om(0pJ zdXHS~FuTqK!6A+}veTu?HUDv(-0xTM@yksnRBOiKsX&esrK+SJ>|0)bc(^sTimmSb zOD%mSzr9a?Rombd%pyxDno}duvG@#^J|;SwRh(X-{q$ zpt^5m7s0fd$>m(jM0`(5It=81og^YkBqRl{d+M)K>NZCA0P%j#0oha4i!60pB`8 zfuisj&h#FA#4wZo?qOpQ7IF_s17Ov=VLUeWVe_%CJ~UNjwo^i5P2Qf3PXhMtjm=nO zF;LoJveMCgjt!hNF##d2Abm#6%5a8$Dh0D^!gWUqq17Lqw8yx#>B3a9T6mjxMBA`} zb7qX1$v`Z*u_aHbVR_rL5T$%f_;aLC=N5po-#oC990S^Tg99q+>PW3`m25fJhsGmN zA2N$*B{rT8MZeqYU;L0Wc7@oU1ug=RficUcopR0OL1!AZTR#(C}ROh!)!r>Fd}B-6-RFOT$i1WNpZ)TQ4XN%f!kti)h0s}9gVq1_RW{&9 zI#*b^xcW`6=5Go)Ek;fxXs_MeIlPNctT7_(0FE?amhVmaTyIuwW+uN+10!p*98U_0 zwiTBC*fuvbKHUIu@TTSkt_QL058D0z3D_4E!2c?c|Fe|1%I25FFsoZ>Eu_idpUi5a z=MZ;TD~jlN8o$-U1t)wc(>YSeFJ5hCDW`&(mw&5DP0l`?joIC26P1_8lAD`rF|;qw z*%_S>v&n5X1J`oq9Y5)|$9RNzER4!2kQUR@{W_|ql@z(FpH@f#u zOYEmoI`nF~&q4)pD4FfvaqzA%HV%%k$l30w0c@tpyRk`3P*BU~{CKANcV`k4(wh^e zw&N>kgN#ixZ7v20Nx1d9kFdksTlZe^^FyV>4kaJj%QyrTr+Tw1YY%z9l;#x`A#td^ zo^^J0k)yv6tCMuS*;GMS!_w@>U)v-b?;UrBa5gROeFu5z@!9V;;=E(x{$Ip6fYi;a zQ`ifl7Q{L2nTrZ0@*U++A3FE9p55~ z|JBJ@RyAXP))F@?i%13a$C{&UEt+Jk-nBTW`E(|k6_$}+%oWg@7){~sr63M|Unfc} z_LPM!&e><9e!UI`fbRS~j$}B3Z~j7C^`_rGaYulk^#t!Jn9nUNbSVf-(c|2pa&m(Z zurZ4+E*mRn!|#{3lzV7Owx-uMly$YLe~qcE(3c?Dbr`9A4&LkdLy+{&W^VhUKMvIc z+6ou-`N~oK7(c&D6IQ(B=H54m-|5vkMTJuecdisS+lIRuA;xx@pk38`w|z3MvJn%E z6m%E+?6OYGnDwSCMW$-kyA)2k&Uia)_Xu>sI2-8~dQ%-c5w*gMM-4!Oee+yX*}q1Q zHnxJ8(n?Mhm1w__p4okV3N5K1OJIH5A2oCO>W0{Za~kDRR0MOq%UX0hwXbfP*Ou&d z9lK6d!ufa2u~MTqf>+dOc6)_45@rW-RhVL)+pf-H`n9}t>53X077JpXq9W5_LlIBI z0s)fV*KPtuDGMuW5A6r7&{Sbp{@Rv2InmXa8t#dV)Xc1kRsio!zL}-Df2oR`z-K$P ztxEtE6C9u-lbyx>`PxeCk>%lh^q5(CWn(p72r;Yr*QLSi9xl8aHw3iGZqZ49O}1+q z`!TDmk^n8L(~Vjf)>O_gF7Y|rPM_G*Ik)wtP)vFFoGKzHD5{}B#C@Z2zuWYjLaF^p zdUpEbw!C!}<$-$F4e6dCieA{ZtnQM_#=t^mW?q5f#)ce7mT6o@uMY)w`empS&l5DU zk?SGh;mFlhYr33C+=4h(jZVitQM2$akiyZv=lj0sZ|^0IO8{f=?GBbrJ}5DkkkPiA zM`Zu2HMwog#~lkkjT`E+&)#iduAHbRM$eVE!1p$0?1-MO%LRb)hudm| zEyf*)PEJlv-t=)fU)`nU8MP9x-C>O_P$*7JOOk#kw`=IXI`sOPw^zf(Dlw!BZplaX0Gi90*)?mTj(1H1viw_$$M?3$$h#HFL2bjY&}t+r^^ z9pB+1Zq3U6CU%Xn8|>5;hXo%oSzjooVySCA(M-ica(f-Th;c3>)vR+{nysGksb}7B zS{om(x-`P!xw|sj7kfogLc)6HFh)tto>_ZfVBjueNqAUMa!coq{<3B0wUGc^B6;b~ zn3gRkdCPs)vGYky35gAof{+$Gg%F#Pma=oj@fmYb5|)yR$5W*KDDzrvJC543qZI|D zDlGph8%P4YnElLZ&V-t!$Bp=GJvEC{h)T~Ar{~|KR7v&iFk#9VAI^oV&AyjB=OCS| zgHuuJGx@7gk!t~?oKYa(4Ai6W;Bi9hPsP~5vJ-|%ds0SZjklY@C7+=8z}_Rb#Xb)YhR9G8$_espS` z&L=YY7=QvtHP7lDP1*~4xtj7dSLcShPmrCBxHR@+Pymv%;Q#r z88SCym6QZb(!YP#R9P?O4-~+Hd1$A4l;mwUP7n=|Xa*x;mDh}jHOHS8h-$D%KtP^b z6S$oVbIDH4lGR4dEP)EQYK1GTE6^jqsj2;gRLI&~cEX|^LA|D6>Dks@6lqo1sy+|&Lh;x8!JEQrL~|N!7^MLBIUhoFB*yKB- z{r>PJaNg3O@|Bb6>$fz6_YRi?Ge2y1aW+@rvh_S$+Sjnp6TrPj#@)&|M7~l3D!q4l z8H^t|757(P?T(zl)7S6qFbzH7tu}Lw_{FkO*d#Xf_&ejXnEsk%(PIR|^ui!zz}j`b zTWcgouVw6q1(sh4%{;yelH_kkhY+i|&n6}uT`$M4+F4}8yx-9+md`Oh{( z3xo5UEM=`_3uN&ok8J|tlUNi<1mATYDGGSf6<^9^q^MU8vfT0qfbqlG#GWPrpR^#*rcTGK`hn*!2qYqh*Dsrk+8*K*m4sn*AP>r<=( z({34mess0KP7n}ylb#!mP`SRn4JfTlrF5~XQ*F$K#SbBwH8tO7#mrBQO-&tczQw&s zWQZxI!fW#~<#cdE*Y+yg4J>Tz4!%)cqq1iZWfu`fC8~Yp+MOT28vrWOEGsV;U3YYa zn+ryVUHs`euNsPwp)I*9@#v8sgzrcI&r8Il$g8kXS*{8yHLCr`D>6zX_ZT0^zr}T% zYaq$_PQ`?s@mkJhwD)Sk>vugXEU^{#-k_>>0RA(vw>DPZK53KVHjbm4>6dVm=pN%4 zOl|o&bDF4{gNM^;b)-i{jIp>6HD;H_4#%TP?WeM`QJ+3ZkZ~KoopfJFHBH}Bu^4;s z@Zoot>>ocqzG_^*%mIhCzDaccK}cLlPV(iPZgU#JFXn2c8_c+PvCHLVncVJ7nE4!` zclqRs>eq#8w>nwT0Bc0fwnt3t*^HG***2fL%tj4^C>HZ`%Y^zx6sDnsykA0vttAUf zrk7*r+zo1DZdLV>5BN`X5ab#J3C!*rd-0+1XTdF)KWL@E(l=xFf{+IlL)i>~QjI}kRr2OAd_sof zV&d04#o(d=>uHN-VcntRIFy`rx~VqaP^7S`W|3RS!g!VK2Q7AMK*A+1i%N4IkDq2n z^v50I<4@xD6d*ipO*-WBp6?5mKlb0^EyROK*TrJ>Mop$oQI@H>Z!~b2ya()y>i1`0tcoX4oj}|99_!!}Q5LSARJ?fGKNeztRhQ571@`lM6oCt_HUS%jxEzZM z(}9irP0bc5ytlYq7={B4bu6L^!3&B4W6W;|f7G9b3B;L}pNRhR)DEHbM|lXreQ7%x z*WGC)SjVSJMvdQ)hW|KA)CaD2#Dj2u?pM3g z2V^B}=je*?#$lr?x%5YQ---+kCx_`ig^I_;)E`YQed!DS$H`l9h92vs4q4sMRBv~7 z*54l`yoZVJ7x)Tjl!mbUP0O#?TY0veVnh)T zi{y0d*5YsVm7?$#OvKMi;P!WFZLQZd32>twPd$W?myUNLqu;ftb0NsUvwy+>X7?yR zBFJ82hQdf*Jauw7TOqRZGBbr2sCMyjfc_=qy&ea~;N1r@{ftBQQ404abl+ph;EAXo ztS>+I&zf8KQ8Zng&3wHI+tvBvK1iPmR75J6p#%awO9TupRC)?T_088cr4bkA3hlTsCQb}^yQvvm;}b(8~ep4 zBw%!{s3+gzQ@o74Y?>6s4?;iUCls(+bmWKh^l;cBD3?}?iY_DjHZK2plArf(68)xO zMDFl1rwLEFw`|p)qS$B)^sfEmOo$6R?r=NQ`i~t?LKgVvPzM2Rv}hfHB;hj(IX@-m z+6M?hh9IMWzvbLQIuw&<{)OJmb@%;cbW42s3rNv_p7sa+aIGNxZo{+Q#-o=R^+$*4 zyY^E`z{~zx+U;Xae!Fy1tuJ^UGx_Me&p2-{G<+FyL4$A8xgVH~6gcBv)gOKRn6=Ou zQ!m8z`0urUvzxYTeY)g0nhrcGY-jt&~W3(h=rC8d+Y ze+&;gmc`)EZDKhs^S6V3Xb^^g-73ah(_{o;;SveKJ;v6|wfFxp+hILXcP6f%BH+|a zSWt#&dUt*4FRLXOW3fSSMiAKXyE;Q5%+~q+C(m(CpU0Kq^Tpg5@dx@!I+`#s@FMQOX!-UWEI>J-Qn0Ep-t&4k5sL?H# zq}iLAyNJ`H-GNdg8M?gSkPwq^?;a1-I-30p!pzOhJ*~E#GH!i)OMf4Olo?W{1@QFz zm*`3IM#`ySPjPkz{L9+lxwiwhvHa1OwY0PxIvC=gDky{j#`xb5=deE8`%jFuc)~`) zJX6;{xw-HR(<`ql5TlD(7CKX5D(IZlt`yaG&oC(@(R&~5l&%dD{*WToGv?My+xzPX zKXaJW`a@-VL5#zZbGJ6mx7Sz@fh@}XCukI8k%In3=ZT2Njyc%_gx_aJj-~^{j^5tA znj3#2_-=-ffQ;NL(oR6%yl~PGKQ9+6Iu63US90>!s6D)c(`OT3tHG^wtu%;;Ne*yU zHsU&efuUMj9-cuge+WDE*iTe?*HaugAXD5{hy}O5Hu5Vw*Bwhb>}0$zb5*;GA#*9LKd8zNZL^4J(gaipgOa0)A1A zvtLEM?h+s_aD46QiA%PvljR5ffeNML+@jj~ySM}dRa49l-F_pA|I+g>;3WS}T2>S# z-R_uI3p^fJe}AAdp5I7JOq_TM`Up`IcFlShShqjpO{rPQFN?oG5eEWFY#i4#i`I(7 zTg(jv(bG2*yD-?zrx5W~5wUBX`|xb(8Z;VS+-$pYJ#q3&c|o7v?}a zW5rY&q2mN}LIcGB$k^|SGV)0U!b$G_;zbC8qJ$YYWp@8EQcww0XDw2}36Hj0pWv|B zMZ~=1`kacxmS69|y*YPr`;UN20OPhjAs+}aQ;JKW(gP@GyhWdia++gm7#t1?0wK4# z{Ds>8@6^UGj!4`FR(0Xbym(1LLJxe4xK4)-efg~dv#Fy#SP*4yJ;K8Y+QvGOMX05S z=hD&9p#d=vMI-t8YmNOZX!c6y<>kGlg;3zPR#i$d1QrO8=%23uaJEiT7>f1lE@U?1VC>CqUDN zpWHER9wX2T3&%TvTsmU{Xz(rW5GiFz0Y6^l)G)sEb-pV6z4hf*;FIHiEoElFi073k zTSZJdZzl@aQHcUb2V^e#{}rB~wRI6t-QaWZnL4T2YV3EK=Q@|D=IQZ1Q!sR&U&KYq z6|T)s#^}P}@pjN)VoYKB_8JSxUrR0Op$g_s=6;26A^~kG2*yCX_UmUD{@lp)U-_60 zx=#^%Bu+!08PKQPj`Q-hOIgX?Mu3ZQ7*^A`B?X1>mbfB0rI=r}gTOGFCk#4iKP`(Q z5<#4OL{6yL=(&%4L+zspL4Rm6#l1kM+V)Eo#)VDb19MgT=&&9(4d75?Ix;#CBkWI< zeRmy~ZMRk%csYnt<5bCDgaj0QfA$*}API@Jn=)TNaJ-KpHjMt{hmNI(WPX)x-~4pxGUp=%!qI%;p;f;%_wqoGkPMMm~>Cpq%1rFY{)NS_)+el|XMlS35ZtGF-?TbLN3iHFJ!%0qK^#r*T z+{k`H{k}j7Ub%TUH~5TmVxANeR;~fD`J?)t3e%n&qHnH4LbG56+Qq-zt>+7{nt(@g z{bilcP|EBjdO=gP8W3lITX4KyuO#HU&L;VWAP)44g{`5jgA5E(Kx(I51>zcCUv|uf zscEx7MT2&3r(>t?VO&ey&H#*T@fvi@DHT<3GZthS+Gpa`7P)W{&%C*v#pTGtfh&6ejWas_ah6MYU71;k zY}ItU$YsgPZDrdLaAMCmSidTXcU-2-B%LzNW#g>|=B@dKQuv)s(Q_hAkN6)(5u&E^ zH%V^aez@Az(GevTLXx=k^LV4VPyFV^j`QOMmyv%e=H9_&9$cI5n4u%Q}(@5XS9^$TsjYF7i$b5s;j-NYAvd?ZR-~)gS7df zNzjY)vr|A-+s)3Oy`ny83GITYPPFIDaoW!67|~@ov~v>WECjG+2wVWJpd+Tgxcx!Z%iK2C=Tt&|K z`v7+$u^tUzPaWo*Dpgg3#*VsK;&x#kM`kyWw^HNczPxZ6%ucs^Fu@Hz&k4;*wgs1KWY3_+}A`2;}kO#d5 z=ZUR>CsLT0GGE`4!O*BaI7O#taW8XRghVHUxelt8vg)b2k3yq3vZ?9VjDd9eIL0dm z1F#PzFK)(YLm0GX?gr%|K)?OFz`uq*1NbN#lNO)`S4prA70JKR%!-_fsCAa-45^cV ze_oIM)>iy=!9X##SxJX;sy%f}d)msHcaXTD5iX{y&Uk9dr@z4fu@U|=?U~#RvUuX; zv#Uj7mN}}{5muG!%OJ0}7U!bbO~&l2Roirz;qqvylV^kyRLFPOuZ+7_FS&a=G<4h< zu5MVk!DRvXXKxLeaEjU&GzA571%PyD&HM-|*RY@2cYTFw@|%Z8?*!y2-+z&IGv>=j zQak1q7q2x=6UpB^+$;H9TOUbjwoU>*E5N6%Mln|=YaO>h_UL4Y*QF)ZbNZ>_42U^E z=k_FGvr)8WQjs1K@+2{COr7;Y^fPDjP;^u}66NMCk0APpe{MW4~jK6lj%ag2|#3EvS6;lAIcLWqJyj4*)0^s0Y1U zz34WWpW0+(B5Q2P{rTk3~kp*XHE?kj`?Wg;h_g^Vzd$tE+Fe z2aJdN&B@5PI>fO?)~aVXz@UxOO-&)Hh0*3SJ~I+v4n#~BmFb`i8dTcH7R*|A5JyFR z+W`q~N9nRrLvwA1bL?@|XO^?)#=FRYFOIv&to0YeEwt;Z6>>|7HVFX&~1O6M7lt4RFy<&4*WBS>F1T)#FXjO7n@sCm?gs@R;w zvC9pu6pvBpI$LAktq#H&*L5#0$|N@aDoN8w*%~)t2Dd% zHGYr-12eo^TmPp9AJ#kljX`i{LAKu5?o%GVWLWbql4WLYT4z}^A)|T-rV?JcvZXJX zSS#?g85_B0u=LLC&&dFF0#RE_s4`9rl28h3k0K~USO{n~qPXXQC2R$V%SnY=+2wwE zyCLWXyPi-db){Pv=3eZ78#J~3CslF9o3yevwC?-jdJZ;YkEFUzGlMBSy0Vr}{rV%f zC`%e1JmOAwtx$gm$Tq-Hd6EOTKR{5rjm(>zP?7uPwU4~}kS=jc|8fVL7~Mo!H2va> zgwPp2r7NxB(e3Ju%j)ung)IZCaYEJ2d*fDH00IbF0}{gKG6-$A7$Iu%ZUHUIr zlbq7UWYY7CwO+i4^Tq!gI9RXWqY(3|`Fy42)6a!-Wuj3>`MRl5vk+e*C`9slIj^Hu!ZhuB%77lWMCn3y3 z%qy}(+Cn^TPt8IaHj~d`ZZKk?w_|u1*y4eX(=gte>qsd9PuKnzF>*^8GuM zzlx3Lm#4~%b#hy&El%q9$F=3bWC9~Ou_L#B(`57H)p-4=4bN`rsrntkn6SGICMg}# zfmv-y*0HAXNkr^2mLxXXs~TrC{{d9mKZRnEEdNFG?2ldWB^9^%P`U6#GFr2JultEQ zX~>4fHom0Q@jA;!X%q$~KPvDzzMixCzlJ(GZ%IBW6Fv~&AmHA}5s?#Og%ZQ)mAJ5VG7Zi@*_c1;1 zPE8|1IY2FGUymDu$KGOGk|jbPWHomz<;=M4Ajj)YWN~%Qny}t*M3vNdL%EQZtAmv@ zHgok=mG#iOp7a{)lcLnHW%6DBP}zj%Wy<6{g!25Ipw}j1VR;eFx0x6uz^;0Wnn|Fp zO~GNiK~0ZNHo9xm!VPqt>pD;Gn{G^@O)FD<9XYgU8%=#E88Hnnr4ePY?s3y1qNK5B z*6LB8-p)?Q`Q2|Fi$~KV?OkD1ck5|dnR1;4zZ`bFX@)vw82{LPV&fDh_WLCqbvbC( z2rwgwEH!Eu5~=!&gzE8X{RagVD{eG7{0dW{qwk|$qk)pjU2D9Ha^#oyFlThZY%=<5 z(UQXpNnhEB8X38aNl)0C-YHIyzTh)jx`nT=DTPgDf()Al7OYnjat@SlWrEG6Si)85 z-3gUh%2cg$7vY;ITKP6jov+d~C3h3S>?;|sdF0`>JUyc%`XPs=UR`vMKNpHVvK<_g3bgaT8 za8<~uqNjH{UrAo|z+6x4E;bnOfim|k?IE_rc46vr$JKN!0MC}U+wl%c0?d2K#X z%$|Fis;YWNsZqU&orK?NxQ7fupg8T}EsgqJ!^BQgLpIb)iLlz6eTHI;RSLy+beE_{ zI@iSCG-K$~qrz(qCK*G7kIR+EV~fV_SIR$I!7>MJJQNr{k&d8ZzkKD|3%Bs3ccx}$ zo|6l~-QDs;So~GUaTlAg7)c)S_?=P8LSpW8pC{p7BJjO{H})yt@0v_hu*;4-k$$?E zL`GQY5V@+mYu2H?05bK#_=L%5QnsPnGo~;oy42&-s9dxZyfB@@z6sD8$*9_r1zg0n zZrcy-dJNNgk^qGP3 zGZT1K2gliXIu_-b_yln9f@KSf-`bFCqI91 zOY=;|#vgU6I;!g!hJ}MOv3^^X702!vaMpwuEKF8JCx`(1gYdEkEjxO*h<;1^$cBNz zfTzY@Wo&q9%(qyl<0x6{lH5Iu0#v1jP!23x9I|0+4DBuou_-&|Mq6W4f8iK;e#?_0 zckX*Gv7^?p1e_z*v0?`0%10il_U-O3TR;$ojg^^%O!^$JIUJ(nlZCGjr+tj~_G_^u zVf(*eZ4hpaM3?6q?tBIhrL=%2ZJohtDh#v-#5`2e8zYWdze4J?2QsLmi8g)X3tkxn zGl-9crdue*s>UUPOc!gqew};leh;sEr}UDTcet2_K}PZL{V6_Y6PJ~_{;i1IW9Z#sLRv-4SSijV-lNwBN&>9hz4Mqmg$qv&#KmVa_11Q@}_)a z25)*4S$<~Ls$0E~X_C}{b+l<5Z+)f~qutp3sbVcw?Z}mX$`0#5pi!>P3XHjj!HT zQz8qWyZr*Omwx@*Ak)Z-Vp$%!@rBG&#UzV?GbW|6`)la~*0G>JZ=mK~uAIVDzyzAr z0p>IptrwpCK8u#E3S3jgdt4U2nq8i^JoAn9hsYAkMReym67H|W@5D-4f$&Z&b^f&1 z)?L>j=FM9M$Ecy(Yr^Ld1$jN~Pc}*hh1Yn=QLfV<*?6Xd_^=Z$)>y7ouI!PWKE2rv z@o$&Jo1UssJ&pR)1M|i$LChfr9fx%R>6VEuMebt?Xmk##;TVs9x`%smbU8@pEYsNI zH&Tf7m|U>>RFD})$;3Bi)t6sZz)wl!Oa<&mD9X3ix3)t%SF$!6{4WIrOv=7Bu;(cC z0_X@_8viQQc}-LQ4r2WXuz^ke6O)wRZvzHNcIBh?G`?(3EVI?N$?B)VtwW7L2IHUP zWv+3acDZGX`SiXpGv;yOXL(A!kx4I8P47GCgav`v7dR)76&?bxt^TF(({Rdelv z#IDgKNQHmM`!IQvEk$Mu@lL>@B>aR&-I-y8`K{UiEaBYnz76O2jOo&mqh$d5_VO_y z^{;7Ar5(YU!|P$4`weWvOpLM$w(=ovFg|DbtxkLeemS#QpguG7Dc`C>@mNliGWFMl z05V4MN-nyO#u;pUB)vpnil?GIB$k0rFZi;lP%H0Wz!va2!HcNqqz@8wX{t{^&wqc_wiOx7@jE>*qvFGt*!>X( zgZmq&bUE_+$!B4q38kv@z2{#+Q93ps4v#wtg(j-`%rvS!>~iZ`pR_5#_a4WD!IdE} zBAfOkHkgfAhow%{EBC4j*NA?Z@>Mpdb*cDP=rn%Hu+vL>-SKXTh)*o@4T%+lhZPOg zONcn;4ey!)dHqDa`^8G(@ulSpPk+Cx7WIIsebJ?fo7jk3zI>V~bn1v>$Lt#(gh$QZ zo_lj(D+Rfvde364R|mpkc$~xtmR1P$Kpl{WG+q+q1C?B&*)ke$tmF7gtmx{+>+-AL z1WI;!;>vKw?tckHc9CPCwz^wYwx%m8w+`c{8&NF$XHVH=m7eYjY`_$du3)!EZ4t=1 zEZc*Dx3uiTxlaeA9e_k_o{YYB0+8ZWooPnJbIO~Vf9%Hj6EsV}t&RtXVlmNB^vmY@ z{AyHShKsx1?_b=9D|3YIOXar@7rNWT++k`q&5~@M>G8|6V2+!8IO&^ptTqWy9s9MR z&KB@Q$=^VRQ&<8Sj?uYN5e#lYU2WdQ06>9*S&2CWz(vePVY2rtR;BV>+>84P;7daM zGDk{dv_uPv+so`~WS()Fgc~u-l3^R{!;?7*5YDK2w1y8N%6t;=Hp|}Jv7EE#bfFg6 z;G5PM>)U?@OnoF4HSba?&0h3}d3s+HBDULG)!d6$Y77ro4u~Lj>f-o^K@j7Pv@&n9u1I-Ozq-RDEYTM2bW8RKns+J zf(#?4ccldUdwk?=hG3`zi{w#8>dh>y9HYy53T;30|if$sY@m)+*md&X5DS*n z_0ff;ZN3mx@n``_h~Q1o>3@_jB{M*gbXYsrI5FeiQdXCrooyev43v;&JMLz?!R}1I zs+jxMWet!;9lbisI|J$)>?kt&9(oID+2fRo8Cf2M@pI_Vrz$`$J}lmarWHvnKm~u! z&G_EX_2unw%B963pRO_Ylv;i3Ju()TzQRCyD}I zKn0z?{shXfr==HzH{g{BqJLqtFM4{%MKw3s!f7kdVhm%)8tn!9oEYAA^v*Mc3dr(VotJVrd8xM=C}xSj}J*F3Wh zt(A-n8NzOe+ow9DpP6T`H2!#RjA&d9??VzuD$Ir-#_p%S$p)8*Y{xBdO2=&xV2|r* zU?+mb<$(2#vHK*B7t<5z%KxGfr=UO=5?j5c4{T%+UHmm@w>f5O5eXdisXQphQp2EC z^d%H4Cai1-(f-j;5sX$~uibQ3;=HbJJ~>1=Ic(tlQDnW# z4_)OFwWU5?zK_JyeJB_G>5R5kxw-;)j~$ot&kPX%C6dC9%-fylBE5Q~!RrjMYBCuL z{4B){kwiglDGA`c&X$^dyKpYYbI>zaN=D3MRb8Akt+9Hjrq+?(q3x02-Ob3@FS>)v zji>@H<;|E_Z@?<3-UdljDjFb?eBvE8qk=d4!&8_uOLT?P0`Qiw#qmf`+Sxo~0nN-_a&2WgmZvloC6eVc4etpA_Eh$$T-oyZJG77u z?LNRTj8G<00hx%zPV0Hnus`H1nE3DfRT(#8KY`flb9=7cvP_Y0=s zv(4Hx&#k_6MI}(dOLFZSQFniR>(Vd^N#ddf1FnPn%L=+o#{ng^_Lw@ndOv+5Zk-AqikLIlix zH6m$JCr@ngKQ&x#Y1V&bi6IcfKYjd5E;6dTxo+@VfYJBSH%`RpT9DGP5~J~A-$?wj zjhk&e`SOld{hp)j)Irj1>slEWfr`-TSYhBI8M;oHP!9T00xp+Ifw>So0Ob%U%&cR> zW9{R@41fNapt`B~LxIBED(YS1Gs_>juDiKG*`r&j(7-2XaDpku1LZ7{@8gPZ>tgqv z$R8fR^J;WRTzz{OpAzn=-O2B2tu_r$Y%?>~QI7_(O7e~KK?$@MRC*EcfJ@Vvyet7e z4nEFJ0PY?ez30cR+?X#nN$(Y3&pXtmR^d|jY$=04iKo?_k^-iHVzlTWE1#f-*mb$c zuU|RQ$?U%IFKzTPxvQMrU~y1!r8*b%5Texol`lXhv;r#HXuFtI1kRNN*G16_CV141 zgAqX+A^@8s?-`+<=>+7WUvjQ5D+K0mO=lS4>+Q>hu_IlZg_!Y23gY#27`v7Zrr}h+ z)k`6k7Iw@%v|8oLNqTp03U@**^2+wk;gJwTB5OW`#^S&e!4RcSZ8x$-gISZmdYswg z&}3!3)0Q4u)o_H5ZcT>_!i)#XLsv~2A(ug)UWCa(GbFZWYo4B6Hi%{d-MWcB%w6tHJ-Ssk(JO=2b z&R3mQlEfJH>`BAwG`~ujYa_dR>A96irWXjvb6*@AKP#5Yj?ZR}F|4@{Of%VmTJ8OS zJ`0dns*guI^~Ky4NaR15z@Tt$Obg4*?J4^73=Y=~3-{ug26$q_QGj*04!cFYSX}f} zUvhU$-!vQ*W@et}Itt>@MwQBq`@kibVu8{<}|n?lN|1B^zt- z&Pb6i6p?Pk3i)4E0Lfr9+m>D35QB894BW9yBi>p5P<2u z05G=uLx)W(kupg_!k4dSr7H{9Curb0Pz=yq$Vx^ zmnq->Lf`#hEGLW%yj#=u!tg))k8>q?4_t9R7RHEZ+(mZ#cqbX3AKEgX4M~l$ZxJ%z z-8Sm*S+z|(?$yX@Agw!=aJTy=v4L)U4(p;vNcz*$Z$$IFX}aC@-c?~lwpGYe$jvI`96*MJnVL$*Npe8gF|-5{=S|RQ6+R^A03Oao?%udWjJymBRx^< z{@RQ3J5B`w2=3n0e)O zr{?6|^5o-df-Y6sBwy5$qSKK3rrvz>FWrKV>=G;H32I$C2E*5X!p;q;I3t=`MK+#% z=e*mdBL}jY);mCbJ^*=$KgzRH+JXqGub83TGK+eabu<4Hf9P^qLOSHoU;-eru+uZY zxWiCuW3OB_1kh&*XZP7R#`YC~rJ1ZE{Cl=jHnI;C?3p^BwIMM1a_oI>dbZQGwQe=HJSTri>x__|z+({iq?4Y~dS z11Gt-r8p16ml-)8OWQnUY$EcC^j`PkaQ9`^x_!HjS$!{FHf{M9&ZweBPB{y(9OVkQ z@Qv&s!{DzJ`Id)3k>r}^pIvQzo0R#_TM6J5A?1GIGrTX?T&=Y?PAh69OR{pFardgh zzE$C)#CMa*rhW>n$9grrNL&qaR|>A0UT~6`@qc^1DXXj16?jWV9guM0ZlG5|aBmMF z1B@>K;QB+CDDGWITfhFiN`Z-gH`GG=#`P#L^*(-GXPq95MZvTr8tAF3zZNXwE%9Xx zJ$~lR>zQk<%k#{^Xv0BNPCEIU${Se)#Q?U4H&@GB9IBUo8v=KJU<58NW_cI{|9%kx zc+ap$NoKZee+6#1!is2}+*!d=nMD`#p|NxAW*YqQ?&?7Xx8=xanXQF*NAOXY)vO{V zlcROTEmh0$T*%(W?@u#1n5VX?rsuo|{WW*mUgHe8yw1O06anBS-~GUJMLSX0GxA#6 z87bL=71_g`km+bb;p=f6V8yN1;VJ^nZFe_i(7zJ&x~ODyNWo zPD*VlwN*$pq%he^YGNpr+{P%EO}T_E%@`)BOcx1-+@?fQ6Eessm&zrBgl!y!$^9~B zCdO@;%lR$sv!Ap7={bL$^@rz~#WU-Dz3=;*HS7I;K2Fvz#1A5qj8uakWSVzxL>guKVDV_c=e%n%L3HhZlwf|!NbOg% z!eDMVhjNPwNWwr6Tg3WxmDW*IG}QnHbZV3KEL**-Nt(A=%-Q8YfB5`rE(5#zTK6Vu zZn;Y3<*=OvNVT)Itw@EN`k`F#HQ>%hPBAE#Lu#{$Z#LMgH z@1uXzO$)lqLI;2p^vSOz8h=w+TOwLS@a%B2=8a6JBnLh#c}dOARve8EOH4O<4tJ}# zLFTE!TmRMK(dz%;VALX$OU;K=F#!Pqk#ekf%k&HDW$Din-<0*px&=SHo_4M?XK{LC zw9-LVT!v@WyBPmH?{2Md44{ZVR0og(lI#fj=@;u;R3 zdHVsgSGcO?_FA{Go#lcw{*lBTxXlIzZ3BW%#ois;x9<-3$SsDwNTR4KUd!g9uoQl- zOcMBj1)y9i=D!BSB3Zd8TG+)8*Fr)8<0F`(M>qYVYvtoZ6HO#m+|$l`V+%poecQHf z)c|xxZl{OxL?qKLee2B z$gMD^@2`ZtH#OU=lQ*}faPp~+Q*pyUIF{>#kop1RdfM9OX97OSe>JQSMsL>AvUGGz z{@!*nIMF8X(kVQs%Wer@^UFB}_pGgw6Z6XFnCtOqYlNvxH447D&=J+SigCQCSGC~M z_}RpHv|_K{ZK-@Ya1Il&%f6B#SPrZ#OS2c_suDlw*8d@PaPbOl;%#Z4spNf5PI8j| zh7?2pJsymzD)ad0mXu%~V943ii;+hNUb`m!9ys|i4uRsWbAN5C2Ct^%NEb=a_@qa@ zO2qGIlO~_vfIr(pB8jKkm9m0xpKuB~@Y1VeIxEqoKSBBe0&@NytSTUzsSXkaZ(9ingC^Ueo_Q7hz}5?X7%pM0S5?PJbK#J&aTGv z7}D5TdSz+f9b*tsRbUlp$XdS0{p=xk%T$2|5|jl53c0AWt30@@DkEfZAV+(5kZfo6 z1S0-S@K0>btB{mBU+nNeIM(`fbZ9>;zBujNk^|S5uRXiU-jGeA6S_~G#%)nXPu@-4 z8}M#A{rK_Y6CAwMv3$OZ&!ji)*9P?ky^IG{B-T33s1A%LdVtU>QOV%3nyAN1S|~YM zRFW;^;~Vqv(pbDB7jmMA+btgcOr-~=Nv`RBN1-^hRC|E6z z4lco9FuOAPS}gSpG7evhlcXb>LU(2%sXfb5RiF+K@@^T-wY$8~dG4zLu!8_#9)`2G zPf7)6Ne_L)zCd%NfYw7b+yZs0`%<4G;cLVapLsoirR$i?z)R8DY!4|Kjc#h^k8#nH zq1co?jbrP8yiuVPq3l7xn1IYk+hfNxz-)JagLwEbB3Y62EHr+*I|s1ZQizNh+jDy{oDxM_RKWrspqmq~Z$cZ&#y@BL+r7I>tTVcs>{@+I=V$Um(+k^z`FW&mA083^Qxz@Ee=38F^=Zc->avn;r8TP%XK!B+*7Kri0} znFP*>Gf%|=vZ$7vl=72)Xj@rm2uK#~y;oo8enSwwmq&VAp`gQ!w^i}*Z+$4XmI_Rz zZgN*{X3j0uX)HGpPqtdjnn-5*l&~BQ>qR`!N?o1zvzrk07NHZYXLdCEjjy_i&k0?R zWA>k`cVVouPlq*R`}FmYa4wBAp zgMf6yw@>uC?(2T;`}yAY_s$>ww z(YjBJ`1sGsJAW5cx^tY4nxM@6`M#XIM%PJ-TWsO}W5+(9jVXL-f0l$e;&ic7q*P!s zL+5oXOA$BD_@84-({3_CZXJmd(|fL?o8q0T$(vbRU;sp?b?02jq5lYhQ4tX_@50=`jG@f_* z{H*~!nR2*?SB)1@c2#t);AL@r;vIqKP)XePh@$tql;XeG( zNLl&r)_id%!}2fHAJ);GJ6BmQ>$8?wvoUB_Y|%I|i8kK)B(uM!wQ04wvSR-I{Sy(_2_?}EyRBc}S3WVADtUUo?cDTRZ5QWk zUnvOOZ?%^+|9skPn3%l$~97o{WoArnM z1@#A2OnR0q1+=biZWaweq*HOG{_8q!0j;run@q8|GxZW~Q*ld09sJj>>57Pmybq|U z3JcsDzi4@Jhh@imq?0Sf?>?1_K>;>$l{%xozEuC}fI`#Z?6P`_%t&HnKs_ry^pyLJ z)cX1?Ej-3-Rn^02I%jEq^O+Cd8{?PG`$Q zZ~BRrQ&p%wVtt^>>sdxr41G6^+F0OxUb*UY)}2$lD{qpX=l7@~*r=*ZM=VEcyr=f# zMXjnfMFaL8^fWAdEos`jKcJ>qUbbs;sA{!+EUEpOZQD%=r)!>t6Z8zL~_u z+miKn(iP+VXRb9A)-5G@jS3ItY)b{~RSZ2&yzcVeb9AY0H=mu86SGiC8?aQnIed!z z@#=#-{}~v;#cgnT?(IdtUyr`(hT-`J2(#NUxcc&}lj~LbWOqX9SaW2OHorf+h}ncG zkAFl(Z)`GTH@N!nsTC+(cRw}BN<#F|7F9^+PzZ}ouYW8Z$Yl4S-gCi_o15FkzR+`F z(8AYL(qmzQjg8G=x2VhgIdcuFTXg?P{_a3%N44nQN~@s5P{77lt0{l~8w6Tu2LHh6 zjwIrG6-i6vcWZe2EA+yue8mMkS*z%Jx;3*%m)Pn4wC=O&Yq2$+F*$k#*TU&@v^IWz z%MESyt!atiP-)v`=KsYmBQsi4a7cEevjNc*u=VR+$;fm{;G$RUIg>6fkq;g<{W4;M zcMN+MK8LeiHc`P(&-ZFatyg8>aD#8PZe}dm^|od;tQ_p}<}|pv^1WfsEuOpE?|YJn zC_B7i_Cnpx9B(;OSE;7W(6*@l=(u`}fQjv>tmjYt zc&ZrO2Y1O%`JDp|bn4KooO9rV+gIMZ)Kr}c;oG?x_Mzm1rw@Npa)FlCo&NbB()vbh zb~^_X&5LwS2WorLN4uhBJJf$Q$%mMXOmSGu+D zpii87y6_-xS?Zv%Y1T#G%c#G3m3T3=;XkQQ2xQ+pfkQ1wkYo%#aG54wE`sotxQc7F0 zP?}+fa~pd5z11;d%^RhO7OK)sKOgs?W<7>#eam;NvcgwXSD!%ZkFm^SR|jqh>d$=ggv=-i6BZ z-`sx&Z_>U}vM8mFcV{YDF!D)T+6h%`{`h*??pfYIW~tfw1wIxv^rVx0 z^j)(YtyAlUaOp9^+UB)$-zx;Zi|E__+-pC9J4HGZ8EXO4Kf(Yn9|)IW!ul_8C)XK@ zJqs~BgpM22w1nHMwUT^D@mc#}j2ht+W(2Nr0&hK?R=>6`DK_#ZDOZj5TT*hb8X{e2 zaM>n9Vfr3{schXhcjlSJ77qa^8K%S692uC}l)t;ddYEA&R?AtHl$_v*HyfzIOP$0w zXdX+y!hcFx1zY=0;U(Ns+*5gc|;^67+vN{U0)pyJoImFYb)Kahl>L1 zB=O)^>UN;K{3D5}79r#4Z4UdtL_FRYnHn)qJh)8Z7rVoKs zl8&Jv26!BhR&bZMq{`uzs)j#(@CjgK6YY_20he z&;E>tWyLN?Ek(wGV77C}b~`lg>cF#$0atT$^CG}4%{AVuZFc*5vSu!OTP~}kr8YgP zr8gxlVa^Ryx@u~ne5Q57KiwK%G2ui9P9~J8m5&r#yp)&%zRp(S%X){s`0JfCgYbIb z1YuTt#&ys(Nu7J!tL;7}V#U);Iy|jqREg@<%rJ$8M5zzn#SsXpiVLbPqH zJb;`ZCY5!J0e0)*4xw)FuG1J%7sAE{u1MrTU+tVr4^#hVYlbtAu+G`CR z!mk@N(bR2lMiwUN^4ao;q*#@n|jbsAF))%Ss5L>TGI*+Y4j$V=VH8@cEA)ZzXW|Tn+jJ zqysr#*nrzNZUhMnJ)-Idu99j~7Ow5b&#xsS7A_A&DaG<%l$QxLn8J(D=KDRFGSgTtq?F`VS z$S!Lr?h9Z8ZS1?z`9$#`LnNZg$hFowZLhEK6# zXzZ*YFX?3MOp!rS$$@NK!^&_YuyRte()1zsB_Py57P?WH7xvoo8CP_%IB1Pe#rv#T z6|VU_sK+7Z6fR+7uAQ8wJ;CSewMPK;Erx7>44j6pY^zq~MYk$9+Pi#HY zEUuR*N0IpvUps2H{-7t&TPg?w`ub5e*@``i$t*>metddq zu{zamA#w&z*1U=kMb3g(A`l2per-c-TrOEFJfJf(c>2<8ik4j9mhr1&w=ZI`s!mv1Y zZQ3{L?6%4pP2JY`%x*H6i>)RJ6Ccdn_WhcEZo}+gcy{pBcDn1Zdpi?Jiz;EaCEcaV zHGshw2xXCMd-rpNF&8z>_!*T2Koqy7Ee=!1wDdjKwINf{3ife;wP4}%epXl$F0Q62 zo5GefJ%c*EPMvg*o#S{KSH1mS=X=oJ8f50n5$&fz_K4X&c9$zs(?Ut>=+o1_IdcUj zjCoIxUuE9!_RqO@a*ghdXJJ(>vzRnZNSeWz5i1&Nd5b!k`D~JB!x&>Oa5a!=o^$YM zId3XmpbGR2FT%pS)KBL4_^Py-eZSz$RdAaYb zjO!vg*0>|eA0&I5#Bw)F^Kn?+$8OCCamz4S+#OJS@!~~BM#kgw7eP3xA#6Sw{rS1~ zHeqAzt4wSC)I`Gio3kt|{HJ8LQuo&IveBl8I|C~lLr&F_C8ece+bM_J7dHtNE@qXx zi`5n_QQVgP0G_-}%U5p=c+L0o`zc=G)e^Z)6J1hY-D=g$$a=3|@)H~%7soiXdoggY zXzk}$(V1!5vu7v7>$j2TBtzX@7W_3|=-&NNQnBFvvt870}ewI33ZTsRd$grs9Td*A;c8?w#}C4du;)enU?vV;LeE;BHf@SwmuW z{i>czqYi7Edx6*o&1qNM{D8%#qL+Kp6{vVu&dq{U7g?Aq{m(}BP{ z<205s+eAC1*g%ZE*l8`+HTaCV>#f6(GT8r6&vfqx!{VWOveZ}lBrxajVYg0c*B+#JGp)MRWf0q}2BV@qk?z&9d*4q@Nt@uO z26n3bw!S|b8Dsi4U*X${ucaF8b#bZ9x`xZMJ8MnctxJOSY5k`l0!6-;won62M+e|;dq#V{I1#>?eX<%agS$30xqQ`Ddd zEZp=pB`o%&mX_9J?dGUxID->9Ix^DiiU4!o9kbL!y$>thw=XtFs>{rdKEsr9A~B3; zL)I7dObf#(6MbcZIKcSPR&eM17gUz8&s-MXZ)2ennQE6dWaHj;gL-j)b1bZZ1gNq$ z-%cs&Dw*Ba#c#)oAbO=OLRM2#b5F6brrX8nAnhVXhOlN(N!`E34V0wDg7T(_zVfYG zkAvaf%zO*I>-!)i#VlJc43&$U_0%BIgd9D=Naj}|As#8K&(E;!VPXIu5#5PmVWGz~ zz+ttTSM<#qUkRVdk|X1VukQ|K?Cghm1W~Iy@*CLd|NQgHI_VXTsTTs~*qsnC}muj=3wo|I?v- z#8UI=i_j5*`}euNoOT(@jzn&+QE$93<)n!ISJNlg$Z~|RnpFMGkzDQDhUzh)7iC9Z zy;*^jjNH7CE{T<$n>&k;YNcEl()Z@+3D%$-(If6CrT?k zNDWKhC&iqw+xYo4lq#m!Z`;+ZV;!@<;1n1+F=3wKpj^$_KovvtbO0}Lid9PKCd5~s z6od!QjME0o zwKnIdzIFFds?*6(Ai}4a@GJ zqr~C*heiE2`t~=w`cL}yw^9Ayj{m0-{y);XJYj(0b2-y>yz4mfov;CIbY zC}O`>0KgR|tx2qdIKRdH%xnhe&&(}GEGAETRa_vJ?xR!0tM6wb2z3Bhv;j@ZCO+m@ zV(jwi+)8Zyo&q2xhv)nAO?deYOJ4&(R0bN(E-O3M4Nnax0gsT2#2~thK0sz6Vz9Ed zT?xCl%t{rb1%WQcEkV9g^i7@jbholGhw5%!sX zJxgJE+r2sIJlnL7S9v_Xlr@P1U>#kGG}{yMAt8M=3D=3|!Nvxe_o_BWecKNi`1CL3 z8J4*gKrw=t4dEY8Vt&}lv3-6(xv}e!>&9kY-Aio`dtQU_+sfY$q-8EP8@<)gJa9fm z#>+@bO6oe|G0==p39!{fhe>b5Q85a=%YFIkRTQhVPepkj9jex6on2B=^4s6ToKxxk zj*ftFYidfC?RH!>5z?P#+ffM^=HagVMon@jF~*jx7GJ-X$S%FiMWt)wx7K}Ivp&Gz z>~((qfscguN``kJh^Yv8^@~Gqt4x%YmEUQ=0mRXQyI^9IKJw5#~R?6cI|S)KjkNI66AM z^LSSvCr1lbIm<`vEcOg{h^jY^2=xXfzql#__WF>af7F zbjVcj*4Dk5;kicT9k&WWM9U^V?2HaZ&DxnT(9(_vXmT8Snd&Fzw1WCfA`QjK>{ce4UB^SjTL1tFY>s%PcwgzdzrV9iD7H{3$p=BTr4Jd` z-vi>Ik0dd!yw!2aXVfy={eUN_NqM>?8-Wpul#-x&`BBJk9SmDj zok4^I8vd1s{+~o5i`@cOqnzf32WEY6*V2Q8D3buLJix)};2_dH60IgDU*%e9(8-c%MMS z%})Hgm921`+NFP-yP)CjTyKtiA*R?8L-J}yaN8xv*8WI{O^=R*m_>)jDm%7ONo=|^ z_&%DvYtp`3#rkJ@5ctRok;V-0iHCEi(0$8L;)eU3{>MqRt|Be?Qy#3D5xDiuvG%x} z7QVt)4J5y(o>_#M(F%|dO>5H%d`pK~T8xksZQdvwJ%h z7D!`GNIOjxXXlZ|fFRHP0;CJ~`fgdNcZP&qE6Eztd?*c9QCo8t-Gu|wGIJeL$KOlc zv@1Psg!0}~(&9RO^l0c-nyc^%(_&5qW;P2IV#T%Cy$YQSqQ4)kc2!L!dN_^=ib6I> zT5GFX2Ev^%>qqh;0v6$95O-mXWo6>U;ZAYV>)e_zXfECYTLi2{DmmsMt^3cY&^OAn zU;6K&yl3P@r&_Ati2yD8GpfB*_3g>F=OVO$gPIR(t5p_<8CRnqE&OZLD9wkwv2oj3 zeiVG_N!C*(H)t2Z!<}z}WhTF#Ds(TwJT%}#Tc)E5F-_R(pe>{fx11O-g1;}=lDFHp zX|FHMXZZ;&wC)Ohr)U{T_v2l#gN9E-Q z02!Ec54k2{un)>;fmVnLnWSsIrTeDRi5M+@!ijmpIn3oTipt`)9W2dKO@ix)h>BVa zSGc(DE|e#DfGmz{J8a-HV5~S`FdbI^wL45A=1N_EJ&}?{Ji7dnE@5}0ltF)Sv{|(> zbv(Zdn%87Qn!?M8h7%XpsvQ9hM#$FT-JBn(_EG^Ja8Uj+I1NznEYb>N%|dUqd5!QF zanh}RXAoG>B8kvC%xsFUgl20&fRUg*=D#_@bvuj|Hh+bU4GaFW=l~nkDYz@$+|fZ( zL35;38B=L)nilnkZ!jf9V#C8P3)Zeh zlsb+K=AOT19|r=ernmk8=Q+R+*x7NQNRgN)9m&b(&pS}?{8NE3L}a|jpVF}KfMCv} zzN5W8ddYvux6wX?MN-#iomf%D|4U`y6LOf zQMy|;q@A?*%Sd8n=hD~xPt(R+^tUO=T2Da?z_3db%i;j$M+sFwU|X>uBwqB`DZMw| z7}9=tEbmxA{4q=_VjVm4Eua}AqzD{<5RN?!!{yd|0Dg9Uz&v*nw zGRR5RI{vYRl^_R*{daSiO|OOWZg&;EZo;O$__a^6EHW1nuIavY2}lew0062IH7Wwb z%tJ16bF&N$Qu5Uc>#vCCQ&)lw_~(*THQO}x?b*jE@MSzd>6hy2iH;Cqb@DVY7$WNo z$cCNdB&~1_PQ18Qss>it>d9wDYAY5MH~P+?!C&QJvq!} z4aWET`EuP2`;Bs0;lEAiD|DET($fS3BoCd>HNzaD?Q z=Odl|-I8G{wM_)E+KZ^P4<{;}|LqBCDX(Dx85!>Oc7la_jVRHn>}w4AQ!>N=mP6<; za$l76%LuAlS^ti@&MBz~YGHlgAVbj^6L1%)n8lP7Lzp?B1V>uh74737j(9|&s6v>j z(Vjjby#y@D+%3_RbIvJp)FADZ3Lym|l9PZ#o}cZ{OAyuvo^-_6-jw(em$Clj$&*ZS zF->5rZ~(Zuak-+|M!31QoAu!!AkjQsA53>?le||O-B;)DYnBS}4EY0_Fic;*iDd85 ztA)NPD@Xlt+ko4^kQ`m}Y1MqUk=mGjkv4OdY#!|3(6=h;T+^qZ*H|cMJ+%Exxz6`M zkgM^MDu&BKuPz&5=7d2PU?5o>TfcbVn<7f8o6k z6HT%QMcef_YEUHLC`vY8Ih}k{AvOF4?a5Di|3w?E-}{qzd_>J&+5~ALB~8tU-U3ro z5C*YXbtXrHQj4?a9)TTt=_K6a#%Pc}Vri)7E~?U(OLX$aRni#92HA!;RQ-6t%PuOa z#i^dm2Ve=dgvH@nP6hlrcC-$}fz}kxt7M|Hcy)?jdEYG~G{nBkJc2tG0%l17qD!Hu zLBwaxdUZ5COvV=HKm?6;i-un%@}IhmOGJ88@}-jl!CygRSB={-Fco?}{cizIZE?a8 zv9ZryKt8hvE0WmuvIT#py$bQh44*`I&rN0GH0Y+wXza{Z4!_ac&meDSPYJ14uY&5M z$q+jZvl+QDl2+PAC7|w*JQ#D|&HP+|3}k{+oq=QJB>fP{(Zk7MYPPApN?$b}p4sP+ zb=0ZUdUzzi&nW-H-LY7cYR@-}V1W=xx3v5!g+(WZdrW{m_w|5fH9$_|?`7>@C?8u) zwx)xSZ;k|@mSrVV;$kF8BApFcNjiBdIa!Xe^;)*ec;gCSEI&UaKK5s{obEf1bC=mT zux3L`INhxe(w0@ENuF=AAu|RPn_onKbJ#5mm@Y%IQ>=WTXz48}^bT7%hNPEV*5ijd zzyL2V5X5R_0p9~!1M@;=F>5(_DTR?cPT*e{v5wZlX}23E`2 z(C!M&pO`)fo*GQAKl%@!DU$UTA0aiki868Qn*5W3#GNO!p^8r6Y+`PKJAl!5V0CGRCh zV}A~ns{^HXv~)Y{4ht}EnKbqq@n~FgUl=TV&v~g2{;YnduVA9<&CP4*b4*O!(mrd> z_5J@A`sWjrKUsv6`*J2>^C33E857 zHF}|SJ83`T_qhjxpFF`@GUc}1(u(Z=!tI2HfijiZuO)c(C?0Wf@o%8K$P(xC?;UzpQHLH+;>jE6h~EDU5ndNWKBUK-}q=h&v^QA!xLhf z8YN=gbgA^3{Xs1tk_zjHpp~xe#}QSkcL8%29prUfrAZ5B+p7d>*@o3AXf9kAg6Nr` zxq@Ke*1P1HQ?bkdG40wChgM-cb1nbupg`>Vp(nE<_z>}~( zO9hC5zYU@o`s*X;MIDT}P*oeqbObKHtb^(A>=d^O9HDk(D&68&?s#DPqe2}(I5oz| z2OI+JU6R&eu7{kAxfJqEV(sywO_qhFQXPVU+bTzpyZjPC-0^R*%8}s)z=lHGs>HV4lOU$a__JCdmks!I)eyN(fI!)fwm-n$2mD?{e;FG1 zDU_Ycqd&aMvfqhp9--Y0MltKxV)kFU=zQuXV~ilQ(;n&mm_Z~M#|69d_G_PD93RZc zw(21e6%l&!G)o}X!U}-OTq>tPEMnEBO1&pwk?tGYWS=o$D}9Gb6kUp!Z{8O?PM6qnlC@I*hqnTU(k8 zM9f)p!OUDR4ciPU;BS=nI%h3IQ7VLcHy~$LM-p5k4zktmqIIQ3q}ZQZU5#DEQ*`2=yQ`3z;Ho}iP2_5m$}bst_pJB0I~^~StAF! zE8fHXfL_0Umfn5vi8CF5{NJJ{)hJ+%zYQ)^0}RQltD?*Jymzi+$(BJ~8o%u-@U0p7 zzCF#1j3a4nrN^9wCF7J;zLMXscj;NHGq5#ReR%UK~mGnuQQq>TvM;82M>67+}&{BZKC$_S`KC@z%4`Q9N)NqkM z$_J)4D}eju>Y+wVaDfO*cED!Ya&NZiK6b@kZUDT(S6;Dxy@F0Ru8cLA=36Em|>zkTybN21$o=p4IHPb^ThIXyqK@x+)bfs%7I~`II3S9J` zK&o+7Q5mn^WkXj($v5v#317Hj@kfwZ6z(m@LQc-sXr8l323g;jRSAxJvWs;+u7IQJ z%Y#JAbjK3-VUH&3q}z|bDU%>ovI^zY5Eom%rY?GouBC!wdu7ndn=p?y-tMWGvq`=o z0m^JTA1yKFYBI|BRfstb>kCx>6tNk4n|DFR?@O&Ca2*I~3Y0cq3g5WT_apyulqiAn ze7%7PoRvt(BAmTu*rKR1NR%XjkwbQ6yiLGsIG=mz=aH>%S@7jq!%7nYtAh1joodv! zu(DD)JlMyA1fQh|CF?TbO#A@RZHR8o&6eD!C~~UNKsa?8}H&0r|)(RACK5O)W-6uvQyN4g5v? zRQ(}-+5nS*t{PIRj5y07XSI86m<%@nLNZYp^g8JkmCddOrlB-vc9PCbNc*0};sSx)&N{B-PI5R*K zA?p7^*hCImfIFDuBI>p2u9KbP!s?8qnDY$OvJn_jmtBfTJhft!eN1(is-rEH#@pl{MBB2huDYZfi4lwzoS21wb*J zYrXg%w`sTR9JZysow}mqB3>;?50zN@BDF6eVv>xBy0R}T z#o^SZ8ns+>dm>%o1c&FOj6sR=$SarD$_5_zfsAv=?s~pS zwd!NuFYrv0`s*MzsuXt3_ww?kC_D=IkOGCOrC^8t==$7~d@DBhs~)>>0_Eh$IntHR zFMgt71k*8c>6NttiS0(IPmcUi00_-ZHH#A4*;j5I&%s2*#Bf4&9H1aC#RHz44NBb4 z9_3)10qG<^AkSZM`JYIp3v#+*k&j7$)}oT%+a$Zk{6bBU14{&ZI&jywJ!7LLHaC~A zB;Y@Jr3m{wqAxHoj2;hB^w5ofB10p$tXtu38gZ@pUBDI=A9M=e*Srle%ty#@2QLbxD z&VscbQO~fRUr>}I!rz5C)K8mO_a}cx>aKmEj zIPUAnPdZ+!AN@6H6LrGr0{c4ScsfbZq#mqLc&X;KmB&vSC+wa&;bg2r@jCQab{3~c zgOp2(cHoOJnpI&vM!w8y@yGN4`~+;I;X_&vYpXE5C=;+cO~ifjkm znK>Q<567xmXVDtX9gjVCFWLQu8sJm*ivsRPn7Tb_2!cNX#LafUGe~K`M}pL&_cw9mhYWe#>#)C1f zAA?^mfT{n~dx9gfx)mQR>p&W_yAcIgu>h@yVw&z3F3Gp|B{Hpk*N;KxNV>fGH`+aPEkO-IhG-BS*uTjD zNzZm|tf#krF=2aGH}qvCepFuJ!`F^HMnK2{rItY@x?oo{DQ*Et9HV`mqUSsj3&s&Ns=nDDKy>DlSt9J@=T4P z6?anrGBs)a^q(_Hf$0IJ`*J**bRocH;lWHA0^T+K!;7=dtgfEjqWjZ$v$Q=~{xeq@tU`^Dz6P5M2Zyp}uQ_cJ z;Pd_@8M^u4Zg845t^Cr9=SDx%0*{d~w=@A0(QC0$I;6KRz>KA|Y)tZeX=l6`l^~2r zaICE8d>jgdJs@bh*AaP#AtZvebGv=&OGc(^GMS#+(}p1C5CrjEA$8j{kh7}xd<|Hz zSF>rN`0;mA+X5`3E#Y|lL<_L!BA$1)uYmnmHX~C+a5V7U$B*=~d*es?_fq{7XWvhU z+#&3H<3`qY>yXvaiNUoGq%S8OM8Gyr;sNvtl5%eNlv(5HNz z$(}v#vCLRAwmOySg<0nQQlo}!-UXZTTJ?Zf&t1#4Pu3B*L#ZA|4yo{m4MBq;I3jT; zp~wJ`xMCSn?%xUr6das}ISXT%>@GQSR~DZ#us?I)Uf=f>DB1&!C`a|`EX0ahHHUjC zexyaxDk5b>F zqB1fBb3Wv>O-AOGwn&U9%*I+Xq^s0u@%u~Mt&5*FRjX&ycSdc>J!Hv&Q1FbuzyZ4B z40Szm?7Ry&t#g@oTrIjiT|koN+ZFg5{3hS5;dO&CLF!17S0UF(qXAT;EM5Ub8;=y& zOFF*ixdW#VJ&u@w9p! z2;a6ihdUCz#`!C{pyom*t-k$|w;*P4by!B$vF}qSh+P(Y)rvV^g{Pr+w=HOhVGUA! zX`$&tlhqL%x`X^SFAhtG8mGP^s`KiGj0n|&bx>Al z*SZ9X8YNQH#K~GD|AKfU4sJORnB5jyUHdhodV30i7wgKRc7JgUrVgKcq~@Qhb!kTX zBs_)g@sf-%XS3vjrxYRCQ|l&ep!x%Xqi*wA-i5cj1z?K&AI~?Io+qOOJKbjHI&bII zs2&BaAPn=G9x7fjdm}AUZ;m|%9*@OvB@Js@og9dNe=}I>UNX}JstQ`DkyNg%F|=i)CAx~Y#Ng20 zoQc*c{-|7S6E!~+!7g+ol!2`$;CVmd5P)cgP{znBL@~jUoMnMS1-##?If6&Ix;wtd6QtT(9;scjZRwX9`DBUF zWUU(uuOwg|iD|~+dHOOkK7%bT@v`@-zvy6+XHR8_J%`HsB+sFqKJf*d+jVd32o#h< zj6^H9>KD#UqO8Vh0uKJI!gWQ+y;!_M<+T&yP_>pz&3fuKauMY6Dhcq z$M^b6Uf+O}ME^(01zZWr7eC@Dtt$++$jA^6wAg3IiemYI<*=m0gKQ_r#C#DF&24HE z5-l1^(=dF?=@^~GFoOVYKU7v_>J~j(ZuQy1&DNIc&b^sI*a>` z+YhjCK6XdSN9GxuHYy3xXgs zN64@r(2?h%dFH_Y2~TQVSC07z+dx*^OthE+DAzn7htq)A1Z4YTOPJjn#BAJ~t65Vq zxrsBgnu#%n`6CQ;eKjh;F{=F~@A`QPQQp^?xlHh%sdG=MT2dBPXHb2vKDZEMune|d z7WwT#2c2WS;BWF(K7;!}EfslGfm;4~gSM<`v|o!V&78ztG4MIH=7@?ON8gk+ zFp-^3(FI*i?ez295FJ28AH@VHe$mW%G4f<|rw36rjrpnD(sKT#rWZGtgyu4+;tG`- z{e|w1K3Fo_?wJ{%@`u5qZIb`Ew$V*e{uSR)+ij6vKspg#gZyvIK4pwZsl@2pl+#A~ z^*E9?|E>gKhMA`$Pwvic>vmECRdDxqS^X&%zl(P&{d_elljMK!f^2G^ zCcTr9R$FHZk7*() zM=O9VEyU~up{&gjC*tkAnP4hv={nwkQp!3bm97Iy`eK2o1(mJQnPDhn9$x?n%-ZtYbW5(qA*+_41-=UH3Rl1;^KbB2!-xlU77UK zx8Xod5Y;x+5`kb#?G?WuIR2Szu3?zk=&Zx*+^D-Y;~B$PJNVmu0mJpr`=D%G?+%~Z z)rh>JQ*e>6^lXV|21XF@w|AUwI{ zR!*@9Qd>D)Yn0a4mQE$z;M}8*gS=w{9@l=k2ph2^&>bBiThgTn;AE<#|SVrtzfmb$o}-6hIzvj zwfSC~3B-w9U&S~vo^vmn4H$-WCR=wTTVvQ5G&BhxG8YqtZ{{E3vw|x=ui4 zo?lC>RBUatr9exli23K&BIo(ciaD=Eh@Ki<5J=5E!_U>tgP)ppGTUq|8Nb&2k~XI0 z*JgSf?V~$A&prnks5igpJm4W;WqqB&QUAJk63Jn+tbSAEcJs@E@*}ebO6`@aIgrc1 zF~a?+#v>8zoTK@{fhqU z9R71i{N=(l<^EoxF*UI)^$Cd)r|;QhI%l7x{Y#J{ptP3Le`l?zRu2@4g61gh*oKyr zaX#0m`l@ick0a~~%bS6*PkZD?`ss+_y3NWB(M5j2H1u9{wN<}H{#5?~9yixj^k2M3 z1K>UQV;cw-%^y=*r_%#?B2685*bW2Gx1?6Jyx+RH&8~bFoCxV}7QHk2hf>7ifDTxsUTcc>@f+pE(m3D+Ek3%nKvIYQ2hPrZklhCl=I zjD8NN3_z%=zS`w}4>}KB??ETfI;{N#rh*_H90g_e;9iUNI7jHj|6b25&vFlZd%+U^ zs#Ty`b?~hN@3T|8fR^N?^HGkn&jvWhUF!j6Rvp%Pj-Ebqa#jHC`VFh%pl8w+HdBUL9($CRPm@|EsY8jx4 zn7+tA9Gu_iDedQy#vxIP9iS}W)k%ggr8bd!b9(g=;KL2kwuN08D~zCU1HoVBG&Y(M zzgWE_BjK}l6{CImzt%ky9X&j3yeTZ9NS{|X_m=@f=4mgRaYoSj=XbC@9WY>LJ$4^t z_bllTYierX1!zyA@=>20c6Rp9MP^N)gTh+VyJn?@7nFI$#$65eoPx@{#XDp~Xs&gf@2hm;3D z$T#PI{J2kRF#*-bXj2jhgtneLJ9a{+!K%e>A`(HeGQKkf?m;16bCM2PteXV|b3&kz z)Xl;4!FN)dfqWJqU$_GL?mS?4Bp4BU^EXgO3KPFdUN#DK^cgpZy+e^CIAH~^zQ_JX8;yF+Q$)mWhC}+r{FLG z8A%@Wvm}BiqPVj}9m6xV_H&#|wwvLeksE|Wm(=Q+KOh^RP>AFn0y!jP3-a3sr zQmG(C^c55o6ru7`Fbms?28@h<)!O_(X94&U1nZ6=bv&q(7d$2P>!H-{&l1QSLImv3 z!7`P_igB84kU17alTm|CpajRNX8TR_#XY%z-J!9T^&MF4Q%EaCFRH*qUxn;DVYagtfw?R-S`AZv3W)k+h%*=%zN?RkazL z^iRe*Os^%P7smDKccQ(QvUBo92HZ9mM=VFcH)ZTLi?}#BTR?l3jZ7d^`(A&z?{r{I z3?!_fMHLgeS$TM{JCfvgl`IDc32Js!0e}M)pfHO&^7(H36vm~?qsagra)+=4^$#L5 zRC~NYc_eK2I_NbJRswyeBV&gLUfz4lO(SF~vgOX-Z-5^7R|{WGvB24(mO6`B0yZK! zSpak)#z1|1wFgl49`pkYS&{h1PJ5FLcFang7NSE2HZ}E~Ux(Xa@+(l?S>-W56_D4R zcevMiDAtG*L!XJ`T_#zXqmVU8vz<+x(cFJv(%2!`hYx|oAzyv2bX2^gZvI-@BPn6PG$^0lCP5q);mSnLKi zrH6Rz;KA}I28++HFTdRsgH?bntx2t$y0jwIDK{h`J&oG^bs=(Ud(fAvKl4s{2GB1~ z+YMlh8j>uJnvZ%Mf1=`leyAl2FyKn!hY z*k=?s=(yzMw_t)k^b_CPpRJPvbtziD^ERN03FwqusxecAvg%xin}LT`Z=vFXHn}AZ zm6*69MQ{gf_Er)l&+CSro%w=xM!ZF9OuTk_XanKV6#W4eG+8nIM{nJHI^O*Ju)Drl zw<*1P`k`V0+rkfqDNG!Yc?bw0yk;e7_jg8c#os>YLL%nuDUz?NJ>0b8ePbix; zEO=An*4}bhU@5hsY0ZP8Qv0t2sc8HV(AD;ui(nF_e#=jwawNY+@b#oxNgo{yH1pA{ zF2Y}LF9RK6*26$3ZHMdJH(-%F6KTw1|3B`;zrK?7PvraUe*6>pek16ABI5sc{2$2o d7lQh2AF^~%|6;1s#zLToP*A&-e*+ct{{dyL*POpO=MLA>P`-4Y?K}+) z%_Wsb4|Hf~eo>^MIW6(qY2XvlhY&5`zf+z%%J*nWdaf@5zx-l*SN$#x4JziszWG_; zcUrecMxHb@49}_mPL11e`O?t%3aC7|`^49DWo-20&APry0`5io#@cT#MV5+Rt_Xjj zkAjMw_lmM}e#Na4Axx=3$|z)@-x}vSsa&$ls+?B?A?&$1Qk>j6ohlI{5T>yT3 zF_uqK-7xF*kiRZk)j8=_P<19LBUSPg@F%mPA)%HAyPBsOZA4PuKG|PCI1t38nV8I> z+E15E{VvjgkT7Lho#zI5GN_C7 zx#V3WynoWI<|!Raw}FFO-Iqg%I^B}rGiJTzIB>(6o6`iN-3?Np;!_uIyT17T_f?RA z`wdlEu2lqlQc2b5q;JWk&{W^XaXMJ$SZ4Eqcb&=e8N`nR5BV)VX$zY7XEQmMRv=Sv zu|J7)$!C%TaP^k6z>~ia`^GF%Ay|Df?}rBT-0HMEKQD#Zl2dk5T_k+}ogq5Ix@EL#xw%-M z6BwAx=6aOM%;C*`(5QM?`?5i%OEt-m)sRwxequkeuqFhxVylfxVg)2q+SDWGz1Gpeh$|aFUUt=?}A=J(H*<(!v8sCo)J!*hk48FzGTF4x* zJbLnGWWIb$?~Q@%5-yTi8|0R-b5+D|J!j3W-3a&+4Lo2^NJNl!iM{?jQYNaw!$RFO zfb3d2A~9!*_~9}?ZKf2Ndg$M4aY~j8L~jazK7ydC>9ZO{Z~!ay`--7|>8F0-B3ho`!|Q-6J;QEU)#6mg>K6#8(M|KCsm^K@)L{|lS4@iJpJ zQ1OS$u(3OJ!&^E7)~xP2J#3WZr0M4&`_$nK&CmCM?Lr7ZOR5NfX{ef5_Rx9vq=D>o zrA}v1D&-Z2r#4~o$r_a4Er&PkQw|1kc^xf!6jqtVMOmNuSRB!1wC)}nGw;~2V~vh1 zTVdm!Vrvw$mJ^Vv#g$dGzdc|R^^SI`jq?-S6EK@Xsp}C~ z&>pBG_}R|TI?d0{0l?61W+?fo9F;fFf_8Sola3y&knS?o@G);oE^2yhclqd?atFD_ z269{4CxKZy=S{Cm(1PilVa_^L7sWc}96!8G*7Dqcxj*hK^#l736nH%Nl zr=`0i*>>{u_SbqYwYcovu(*W2#}F72Go-ycNDy1zVwoLtWR=>RIrs26do;<-lAeps z*l&l8m!lD^--sXI0xh8f4uit>*4rZB!Ha1JoY{l9cBfwCQXh6MLkY*v-q@1Y3pp@! z$i5;I!oV?i%~chgQoLkye^6O`Nc;I3%Rng7{&Ju9lTxyGb zcQE9%Jm&bp>2>jtR#;~JbihE{&VF*&T!K@lbHgEMshFOThOAm)D$rPIhq8^|yd`(A z<}&&ir8#ks){S^@fUg}F9gh!`w8W!}y45~T?J-n;;264#gSHqTei9=$&(NA6az->G zCsHl(1JZb+nzZ+k+sL=NVK*{4_R_ExrF%PPO`B_i=z2Rj(dGNn$DM><*cMvWv=20X zrI`Va$Uwi|DV(!0NN4IE+x~)zxk*2OM>`t`8d^J zw`xCN$U*va0oU6VjkW>uF{v2e2DnLl_plpn@noIy4?!7!2TRxFsrgDMc4ni(FBtDy z*^N> zb$3`~0#+$AE?GB>Q3mYwXEb&BoO^F>v^@(tJh<8E8>R_y@CS{x-x0)%?|yh^eV=UN zbV)|#G3-kl&)7ab$X7d9_9a@yxtbO>E~8q8nyv^aEVhvuCC+g;%8n9U4>&6Z)2hnM zeTQmu47m47uL@7+s5b{SP6Z4^=+gY;qPlp=k<2yls*n%E=u-i7((SsIn<=$_@-Ta4 zeNJ1kV4X`afNY|YhK7kTioTt+U6q@uT+Pto&z|L z6bk{L)HVXX+4`mqvD^KWh6bGBr*ISai!fEtGskDvv5x;&Zs5;%Q}c}*GJwxbQ)*Id za6W>$zH|1u?>4b>_d3;U_;Bn@`X-LfZY|HRAK&^jpwFM|zSQN}$vt%;wcbAV3B>_T z@HQyt@2@a>x$eV^Z=>bH862IZJ=$%MYU%Rz%LtScfRu=`5!?zcDRG5`W3ndC@LVY& zmSW9Wu{DV`c-@>BR3EaBOvkRuy zc*CKr5RlXSg^ZstTcdu^#F1`c^C2Tz+`@i_@;NMQmJaZe^<0w18i%6uCu0`{{Bp03&Jw`xKRF z0JV`Tv3`SUXY6mZH_uo$vO)Zu1LX~FYFjFo&{^_+^gzT1WFUD9@5keWDh&$_JzOQ~ z49m&Gm(2L5@AH`tjK(Q_@Uo+5aDN<*{nvb{_4eg1mbh*{sc`i#fUf`YopFNuqV~K* z%?`@+VrJrFIhB4=Xv5`nlfyuH1hPOces;p_?Sk3+w*!hoP@7osr4_wMabdYwKM`80 zLSVE3Ux=dtR3re{;xA@=uJnW$lMFdW=CIiVl&5CLp1Z*eLHqVBSYRji{uNWHb#42G zlBSIU&iDnP8ZJGb-iQny+RX9s5<8W-3ht=1;)i32FDdx?d^%ddrQoBkaPKHq_!LRf zApF_BQXpo6PIt&bCi(Qkk5lOIGQC^l=A2&>c`We&~45c+b@do9)ZT8~OOd;1lQSrww1J1nvbUs&_vsuzW7eh|V3qzm7bG0P= zDqbh#7>*C6B=5(AXz#11xvM@d@Ytm%?fDtyWPjD;GtuMwf!4~duYkcWSnp1Wd#$#p z7^TdkYcVkiFw}ZN1lKncmPBP)M`wQm(xa>%%Uk!blHN-6|#hBCm zw<9?*(d?fU%9GV$Xf#8_UDV!l^Ak!(Gi*c^+o(!tiiHvg>8oW^*1Q zG22dlTw~x9{M}+2ca#5~O5;^FwS@2>^Il%_X2xO!eNyosoRjMr^&SbfHj{Wm_sCXM z;KoXhS*kB>l;)j9ak;{AP5~kJwH0t^XT$nB$O`>SxO|$x722Id$+w!njPI??PVU>5P41*wNQlKx7#?M)A{sv>o#(3`PrF|rQGu!da8uIrwI>;Cp zy)l=CohkAUXlBtrXZ-$?LNN}urv)1fxDeyutvyuSN}wMJHUW;L!V|5G&DFwg-3xyd zCx{1R2^t@o2LwT8q5WQcc(cUk|4fZ}NkIC6$$c&mEjBtn_)iWLZ03ZSu zXrrVGMiNjO;Ib6x6z^T6Ps^^{_{x2ey)!O3c;P|X@7u5aBNqAmJ>U0wgq8PTV06D(L0mi?r<3Azoq=Tjq*2gFuvMSAhkgSn_Nqub(;ACNk5ELa zB!{$&Z7=xZ^knvR7hwOc2T-58Q*27;eyxDmM_c-!R~9?Irr9YEKqIHRzrd4xe`ShP&t&m8NUI#J*G1`XISj4CM6NYJT2~jv zIH8*eln`x{2J$P`rcvwuwL&h-q z^52SKzWQZsXz>U4A%UMMpGeXGSeXSJ#`=Q>>8>y%Ysyrx(VhL;*sAB3D|haKANEfGQS_I`bEu4RW%>p*TJsP z{Uu#ddz=l5XE0S-Ygv(*7E0j$SgHD%?B)l%ttix+;_^I~g}eLw-WRlTwq$Oz_*93C z?QEgM1z_%$4FCsY-y5*D<9g*zZvob2p?4~gT~qs7+`v5`9F_qbfxEr0#!!e=F~3u% zE{EJejRRZJEd3Z-Zu5sQUiq1{V|*F_klKKuXE}DgXbz?RS4Y8rJ~hBv<;^jLL|i4X zkP{-oT1D$VM<>xyj$O@o351w_K5>0x9kO1-o;V6>tB?J;;@MF}DeT|S z#U&R=NjsOlP=OV;br@KV5-Kc;oadQt$*3dRD^-Ai?-z`VX;o)_THM2NlX%&oO|GJ@ z)`&Aw2tJ~lpDjGfI*HC!2*cP;x-I=>a^k*g(z2_0V%p(tqVADdH(a4U(WTLkd8wyv!5;z9vK)JXX)PQPg-2!ygpq z+Up_h7*eS0rq{$@WUWJgODoozhTudK8e^g5eg$4}?adR}v4mnAU zyEmk<5jcoCN0mZRK8I6($(A;=NZ8L92&fgQHMw3m$q;O!5<7J}v)$f~(w7g_+%QW;7a#+{0!P4}-j>J=` zP8-u7#Dfi$b+MjK{`^hF%}FKHm1L@9op#LuX?YTdjgW}WIui&S0K3Zxlz1D{jnge4 zgiOsGrDQK>XT8X|sv8H)p3lHCnTI_fC;T!^*a^JwsL7QlB`@wZDt#y3UB*8&1GUXv zyZM)Zd^k5;nVrb{K6GZL|q#Q?*og>Vf4sNPD!GJMRWkzK4QS!1jbWWWtO*3 zl}$NE82xo`(2koufET<#^DGBAl@pIN)b?G#`$j-fQ zu_z#uV6uy2>j9+Okx_itTx>Y>6Pa2As9GDht#hMuaB#d_4t zj1eF6y)^RbS`1q!B7?$L%92k!#9L$iKo^7C?+h^v#o$_6mk{$#I_#p)`X(b#oF%Pc z!vs9)hr-;T)jj$u^_yIdg#>FM$aQOPZDyBV)7$!2+nl$;=enAxGSzq7UviM1s?ABC zb(A&EYqYO4e?Al8>W@YTKO3$0>Y&H^Ry!BLW+(96;MJ35lwlFEm{94i!^oL-Iude1 zntYXx%FA~F<7NE)%5ML-{Pv8&9%awI&(hcZw?wMt3N1vnQ!UpVF@il))5wC*1j1Mq zfyqon2k#^z3Gp=EDzazBvPD3K{g&BydSFo&HdLWt{?_)D5=XXjBpuuxX&RKXKhrts;pi~YNp3?;b?DZF|o z%k;2~`DW67d=PUO|D^u5gW0ILcmJ*ObjyBuKUha{Lq}OoH5)E$(w5XIghp*4#oiNJ zhVNa%mGQpgsU)M0nuVhb*=oS0>64sa=cO)CgWP25p%;>dmX7ao9)Hq?_%MqcGJni- z@W?SDJaax#j;9ViF>3nZVPwppaspopi45O?nmlEYRVHi2OQ)wP`3K*_&U-8=grCS?|k|Yv-i0z^V^p$&k~nJh6Le1 zM)^gVMGn0znavp`4Xh8v7R6OkheLahZ@hRR;Kse0{oUtaGrdCoP7LO9FXrZzd_Vp2E^YwH6f)QU z+r2ZZU;=B`bI?KhL8td*r_z-RT+qoOJ^Sa#ZiJT{<36C%#r5?G|_+>ya6K_ z=`7^u#-hoQ9bK~Rkm>kg%9hoiWiT@=&dWk`i$vYa^)OrnKPSoq>A|t z2{MoQRuVdCw!5v?l<>axyg~MM$71LES~;~$gR3I9FJMJJ+G8p9JstBTaVjx$$_7-t z@o)yk^QyzAWg^1TFKf)+Li)OK!d<~mlh}Hw{z}q-Cu^ADJ=dK!R8qYe<@)!BCLfTH zWi1(ziyJ=^-}H7S8iZMfukN&<5;h_Y> z4@!LnRv7>ECbu;M)FP57_0I`4i`W1hI+H8h030#Uz<$hN=sP3^z57&X>h~oKKWEdUAy4N z9D84FigMMj+sV5Q+~&=^_E07FRyClsAh-_TOs)dZcvIhP@qOBbQoCNGw0kWxgLL{y z7JkN_E}OdB+!v+%S_LHXJ7ysRYxCA0;WrIZ-Skr3aQ6~6hd4P!<7XqW12rss*6wMz zr9-ejXKRyzVDKNFPfa){iz*h2F&*ujKw!v*2M?$}U8Ne)7JWl?_niolM#DEtcrpKd zc^8_sZ!XdtJw2-YzqaHkk0@<$X%>V<3$hPrFX|Sf=>}(}rOmtG4Ix*X6EyX2(N9mF z@rH~7DQKsR6Mm5`oQ8kYrS=#xw52G9dJWQz=k2|gQ$r;k;^){p7&C)S;Tem*3fixs z(m4UKzY$jWc6IE8^?0{Km6G3F(JT)J+f|;TY3sI6s-2wU4`OI32P*A$O9 zxcdZNWEC|xRP9|A-D$!O&*uk?bUDhFc(Eq>w*#}frH)m@a{&iOa0M`@qbwruCjw&O zN!Rl<9iu$n``$f!rz}*b?qttk%H;P5DIO0bQq}PS9|%+T?3q z^Mz78+hKjIr^DI90IvA$6tN2)kiaTu6T~`d2YnAtIZkh?0!O%6Ca^$tL8mjrSU3++Y>lk(WmvL{lsnVgiiIc-qObWxQu1sHucp-@p*-nVok zYR$AaE6kQ;zV}`Q)^0@Fexy^*ZBTDkE{o*%?t4D+BSmRXJPMZRQGE&6@M{1Rof5&k zMP6LYskcf4ZP7-IP5UKyFp!>|lRRhaFm>(m0Jn?^>(WGPnQmYa%!W2@aemgQ+V;uN zYbBCf`tplOE7(At536UjVpN_HmaEcg>yWiB7sF3~0t$^mYY!3?4sQl8o7~@1FD4b` z@oNv~$Rg{pc%yLNTB3sO-nS-JupqlS>(aNI<(yMS?G9Z{moFWxpCY^-8UbqF_)B3l=?Ae@o<&%p=|w$bZB19q!#!SpYMO5h=$yD#c@{O1y4<5 zXE{37iynB^_NJQtDJisk4#%0$m%@O3D>RXkfNT{8=B~_40x`C~t{*`FP-mZ8VUdL(mS{lN>zaaSPw(D)EY zZJ3>;&&b377o`1f=LPwXUIs`>4`h^Th^xqVQB4yEXa1zl*2j!l|GL92n4n2h8}{h) zd~ow*CB@V!NIK+bhl8m4stPACv4qrU(*(H|42_hxpXTK#;S@4^=ZqYVKvkw1ME#Hq z;vMCcNH$>Duk{-f&Gpw{^Yhf6GVBZ~z0EE}HxQWCg`G_`x(&CNPbuEl;2!QA`2`k2 z-(X*c34Q8M>2#_|gXV#wMh0~eY~@6Ow(0JoDqBSOa$6!+;LI_)=myXQq+%@`vhEk^ zzj_UftZxe8AA~5lC?Jj2hmKUjM;~$}WX`wD2@KV}ldy|MhL`~1NbN~D;y7lQj8jW( zK4f>V4vj@gV@%~cVQ&r6Yw#w3=g%99y6MIFdiXs~^wSJ6$&2i!9P|SVRABhUNy?Gr zMxXT7!b}HDW}6RlQ4P0AJcq3p#?3C>1}HmN^tR7yQ{wpbTjpU2@Xk^ zulT;JFei1v^x7!NyEkF!+|%5(HpT~?bO*G{q$^b3Q(+}l2{YcCt#A~zo#`Y?^hK4^F%C zgyM74!X-I@NCQSa%bk1s0`HgGz7K^QK<`c2G3A2!wahC5`h$fPVG4E&G~`}0c*wp6WAea zo)J4(Tg408enn@K%Tc@3IM+Nw*HZL8Kgns1!8TnVK|230#*<5(paIOv`RP`aqCD#b%Fm^C}qQDG0nI@wWhceph5W_cUg0|SY^-S#$ zy81`n^&1@7duD;WOI*V=->ABOpcF~J)9CtsJMbs&dUDoTiSJ};FbT-i9+g7MnTK!< ztZmI}CqGgAX>=@<$va;rfYWKdBSrfub^P=i6wU{@z)lnE!e$h3L14EaQNwmR|ZjxWig{wP(<1A zJH6<(Gh`X6`fpCyhfj3xt@RqC=3kSqlBNv17i$+8@h6DPs#0qrCd~+jtIo(7R*s6L zL$DPhT8|8hBVIx*|i&D>9|VvZ7n$&pwtnkr=`bj;#BG6yO1E^+}HxCM#8D zES>=Ub4L~DAO$+6&<*d zY|0tNTQ+0vqI?J!kWmoSxC4I640S*MB18+2%2gEF+*p|86XID9Pb7CYWjUlD+S?QT zx?fe5-`dP=@LKjmZIgD;6>aNNmqi!d#tbhQSjG`5T#CBqnPeU!x|@eW1A^NPgVu%S z2A3u036MJecIMRfwg_Dp>4%N=_je6dU>AIz;FOnd0g?d~0!v(GC9Zu%w&!r(YfYgJ zt*RxPUC0j}jY?Af&g`M7nSV_eun7RyV;ng33<2l#I4&Yni&2gd)yItNdUD{-z}gD+ zAT9LMy$+Z`piCr!{v`82J>dR}3c$O)9zJy`%)zJGfS~aOm)3C(IQlrVJ(6&eK|a3k z=aS>Yw+4_AxX_y#UB)nrY-?Ay5%rExnpkJ>m5H^)oQUcax{68fx*yMn@}-4cROPQ8 zC&8aVsJi;ePCN8jMa0;a!(F{MOV{=&SvjLEu4yr~exMXl?>nkV1L?!!L*~~B1v%4r zmt*ITrc{WU2M4kv1_mYbmc;LGBSRnPT0TJYhUn#v(^-<#|0NyuyzhK{$V4B2hHud^-0);Uz5lH|KJ%okhzj8T=InerUYGw(1%g@x^WJ5ZLs7iXTP|8Q zNbpr)#G`0nh0XNJCAnZR6AK>lCVGX}aR7Lqp9YptXPhd39B@TF=;EG7T}sdJ=%iaw z70dBZWbkoh%?ph;V%leZQEG^jROOcj-rQ(zN7N9+jjZo0rb4LU5i3Guj>Z`g!KcB(f`0D!jMaVq>xBF7TIb-IbGBYrKzZ*~|LDqO30M zGfRvBXxi?*-S8pvnG4o|3)Bjyvft2j>J#DIL?{E5NgmH0z$3Q2x?=X~DCiw*h?6Z@^cQUMI}uJHk!m|g~bny-`-O7wgU-Ws!WET3!sT; z3%&xO1%E6|??*-@R_Xf39E}vY{wJwFB>_h; zrxvJ_sRe+T=@oN$h)_V~MN`^586$lTskMF78A^P)97@FxmZX~m7TKxDVZ{oQd6XjG zRhLSa(mDU$f0ca+BfZa2zq<261d1#Je1rd3nMAF-u@FzE1tr0+s8nQLCKPBH)O-y$tyl8=Y1>5j)pizkeD4iX~}$RgWH-` zzHS-huVev+J>dgTgBJa6a#2!VnE6Gu({&#o%X!Z}#{*q2i_<;P8%n;aQ3V3;G#qN5 zrcvxdfU!OSMgwR@$WE5Hp?`BVzt|cvz9d(hox9W=UB0ikGy2{ep4+4X%hBveT)?R| zEwqt-CL@z?4U0T!ir(tGCxPgZLL)+ZaZf~0_LZA`@r;+=<#|qTbhe<%ubO6zyq!nq zY{CtEsZUi5Jk`EIc}h%&(UEGi2yL{HQ&(28-Nxf5i;H^Uiy#CBF&kbjX+$cz-_JLz zwq?s>{^+_)nq#}qW2EJIQ;QJ)bOtmQaHu{?cPb>8%_wyWN&bn9g)JFR1x*8!{nHrO z0C`H09rWR}C5B5*>OFBzvpLu|3hf0%I`t3=!-0_sD?`5>u9)N{to>NgW9%MJQbrSx zIXAH*SJ+jMQ!-+XAeDSJq{cwg>%1!)3n27ye^m&w+(cug?llFmqg)sXSM*!aG|PYx z%zs^zai(l9HDiG`nS0&bcMg=4BAGngrJUlcWhm61F`61cQ$h|jB6&Wg4e%OO6!Akh z>AGVr{i%|)PC+;| zLT_%;21d9;h)sK4U=t?e@EFy-i~8W~t28e{fE^ZrtMBUz@!qo)5rZx>&f2#3>v9Pm z(V=*qz-Ey5DNMiUux!rtw+$qj8hp# zSEf2+CBva1B*&}Ux<@wHkyKN0I~YRDHt3+#ZX>60ssh#yEZtJ++2LOKtxnq;K>#U- z-ZULMP4jc&*qfOx(TuSP*6uA^<&f$&Qt4&jVJWwVj!dRq(zSI?pCsG>5{~x%L#G4{ zqTA5SZta~uOGXlT&&$AGM9%(Wh^{q>p?ejJsBrG&=O? z_y)ZiH#RN*F>E23n#=5<^2pV~&+|^`Wa?agMs`i^vVZs1g4xKyoyM@r?*>66pD&EV zr8S1`0W#{?gf(O}K+KujfgGV%>r?b;?_6$6m43kp;%fsgKpnGyT|-6Ze!dnG`;!ae zZG%Q>Pg{h)(dqwcNk!+eNqpvc?GsA3Bqlm|uqqHSwWg~zW$T4y>jAS~n^bU~9os0Z zscoxvd+ebnyk*UXQL}7&v9K}=bWurC7CDcS4K|F}3}f6A7xl{dK#Cn0K3rZq$vcs_3H`%aPUudpB!h0Au0 zIQtsfpcTHu*Wse$v}0f+nz*;o*p~L)X=75{csICE#AXmVWY900up-#oS+V0Kv|aP6*7- z+k=0EwdDlY|3oz$xS{b09*ri=fgh2 zJmPsH%g#_@l@`2INm6bsB;)lumI1(TW z$rsMshl~X#sjwGzLweOMAAtCYRLe@OPQ*V1suS^D;^&s=IeDWQVjefy>C6Mgi3q6D z!6y~bvWyLOENBJ-rhZP`L8&_byq2nM^n`>t~mqubeU7L zCcebCWAwU{QKxu-<7OK7#go9Pt4^jf`!NKm$P!Ym<^4xmL62QlBeA;lJ{JA><4(xO zJPx^KLD$VJh^ZI`xI$ZzFyH;Ksx+wMJMFfd2t~fcSh1lU0e|1{MtZb_|}8CmXR@hjG4L*%RZNVI!9^mf!qH`MXgX zL(PW`pw<#UBOEo7>rX_%-@5xBN+u6uO#+$IQA1#*G*ndoR?Sm{NzzOatEplkhW{Z^ zHEH?BNUo-cS)NS-T5J%e(37IBJGV#0AyDlYy=L5xU}s0D_jSVYse9CFFXIH zXzqRI>airvn`R^N&-qW{_+a4HKpdZO{)+=s-QwV@g5L)3Y?D;ft0E^9xx9$}vQL+7 z;wAI%Bh8DjvO{|6ZRoL?1ZtW6@xcoahHzvF&`3L1afRhVcuwnE8}-8ijRua^rR2Yf z^fH47e(pesbK6E_TKRi;Z7x#NMt!?ab#V>?wCsxYoSffvC&e3)f-rJ{n%N^oy;{TG zvULvYzX1og-YHkSL-70qo6rXn;gRa?=G|glTlcA@7{_q~LKwQtL@&fy)E+XXq}1%Z z$+s!6)TY60vZT(@IJXwZK!Q{$#SP~U8{XE^i2YnsJ^gD zoMiF1dDH`m<`brHOnWdOYMTGJuIgk#|GnivnOd9!oY?V$Zi23IwN}a;wzz(4q zj@HFC0ptZNu1sK59q2MNqzos1UDR+?7mY}ggpW+MQG1o@tMcfb?5fXg`gJz}bM*%` zpM2!Ejn{%!0=+~b%{>tx5%Yl!jqMt)t5NM1ik};!{dSwU4ck|jcPd%7f9KrS)iu4j z_|0S1mo`w8(6oTtWO~#pSQrdZ$}~-&99?^Q(zld2&~b6CU$ZLtz6vX8yWlJrvq*bA z{7o;qf8`}(J_2Zq0xA?}6o5F)!VFp7e6dW0-wQowq!j14ovYR*W|`#ZXSK2r!@~o! z4G!x2n2W*w9beh}!80E?_TzKn7=Hd< zsPRfOr(UDa9Kj!OtyDDXP&grU$ow{3f4X2CUD5 zPi*WhYCD#HB^&8IW8U9qZi=P8&qDoBK77@nY;Y#&V6oL9jJWHeJ;mDhAb=ewbS;@w zG%(--by+&a&$*|x&rP=K=%IEy`N1gPJ!e_>fxC7aNe9K;;i9%ZEeTTDLXBj*l96`# zSK*6QRymjc*7J0nk5L%sw2?#IhIal;xauEgU8_xWxvua@lQEewS@LzI8jQJ8y@~UD97$jSAgUIk&Ey(>?b_%PZBnI0tI(%*p(Rj$xGR|X zAw;*LIS1dNrmG<-nbGsEO^l_jPscah;#L|isl(WG5;!kDR{^+hRQLE8hND^=XG zHYka>%8Zsf+S$1VJUKgohrw{ZsXnVgfc+uj#d;A5Z*I8%@Gaq)kZ6b&@C`p$3{~x- z1G*kJS^Hy8rqow3D%Mz057+1Lq zdlV!O7EhqX<4jQ1&@k$O#Ku>z3kh}+=M&$&?-D|k)#v+u<4-~kJ2CIint}!e>g|EP zoRhb#906}xk*ipvTM8NHDi(!q&Qw5x!VimJ*Xk$A7e@;j=_BDx7o_64wpf@%^uM+Y zmTFIpFXU3ymD(%5Pf`$pzwp;dwi+zUuxnS4REu6D0r-S9^2k?qQ5@lc7;3XG0f4|g zwi2L{@Q!}{u3JYki%5hj;WF$`@I|vKs=G%6^b{S#pho>9_WliCg%|++#Xh&A^aJR} z|HGHj&Rg@u*$hQ+*G*;y?mT)akgQ9L)lqhO)K4XZB90Q z;npN4DQRRpaO-t3;Bf0@7D+oNvsg{!_IlcT6HS-AXWbqUVPgxrcz^i^Sbg~YP(rgl z2gV}ZL;ul;KzE?6Ir*D6rnLeEe~efCO`)nL5Ix8mA@QJ*O}~m&ka~-Xpm7{x*ul~Sh*?GDUx;C-b!d^=xX77Md6h~F_Z5yEQ)-dj z@*7*+g-B*0ZYLdec^4T?)i->@jxY~Z z&`<$Fh(Ql{3avOOvCu~tncg+71ZYO6F_7d6d>Siax4vwu<7)~wurFJ@LLQW8Xzy<&sA+0&n3=P4<++0|0Lt8&@CiEW)hS#EIox5v-=H_F?2d)w~Ke*1u zPG%tzdCVhBjL3RMAJ~D;jQCy3D|?~i21=j_z#_4GJjt4wA^2i7$eAj_&j51 zX7Xo%FnG&Kx*}LD4>Mw3zjgi#tOw|S%u8tjOhBV10J&Bo;UAnnekl83#BklO!{S_k zx6uCqKS0F7SkTxvVxM{{x~dVD5HTZJIBz!7+x~K?_r8;{p=CppdWU{k$C)p~qjxbM zz=@3l%oiO<_p$}JZT=lWsPDd4Y@CAI^+a5l@>FSBNCw*74uSsVW9AEx5yMwrIF;AD zprt#UnUdLmwL!k8wn1j;1EqvzDVrVSALhm(I#T7eC7B^*@$#fpI{tlN)&{tzc>MfX{0d;bwqnj1vg z<_D;V;Rnu3r1^&`dBGXyDy`Df`h+80=os6=&N9b1QQpH6f9F!`R8XsKXjfQfDd3c^ z?SLTjd7@er_cucR9YSDVb!>Y#;-7b|r_WOo)~X&TI#zO-hCKORjm-u^1T$(!3bl;S zmK`YLv%Z79o(J*$Agii)n-<%;S)a=#-3iFd2K-v;%uIbPpNNpc-6UYy-v^*m^TOUD*K5Am(fvAn5I!s+!?!v9@{PDn||Iyxg zMm3r4?SA&2Vuuk0Y`~1tyGn0S5fMUBP#^(eLJw6Dfj}U0R6v?2B|v~MAiWqMG%1N1 z0V$DzKp-L@ARq)4g7kLoCpe=sd*1i#IcvRVoe%q~E6apDd7k^e{?~Q=ZZZ2bB8)yh zuPB(I=UQB8dOs;yjIu77b*pX3{~AUvG-fSOk8X9%jBZs%3oxWypi7rn-z#z5Nvmg> zd!j9*Uc9^&d)UDK^vGo0fD-cd@Q7x5j~%P8$Ae544vZ_tV!u~*nGaGRgw6t1z&0g+ z&~Ah_1R$5(7WVv7sf5n*KYjU;#SX-bg3nbo*QV@ECtgy7^zE)YoovJ{|;Qexs z4=ro{Qd%-@oCto6b|7Hzf7_7+bA_;sR)`Gt+I(w?*kVROLLC&>|Eti#QfF49PX9&# zWeQe3f^87uX@0pz-CSPpy0kZmk*V^^i!O0r!Q($kh5v7o3aL?e_VQbiGHS_EN+Mn2 zeFtHB!)4RwWi5GzEvgAr!|Nf?q6AkQ4UHn&IbalVE4!tX8v_VOyk^4!E)6Ho-ZVgsNE3} z;@qZss+{#3(4{^Bn&*H&lRHIS;4CPLMdeT~^XF~n~Sn-oMv4Ck~@ zBm#!T))x+@5MSAcUNEZM7t9al>OX(rX5qeFwI}(=Von=zhpZItctl`2v>{!JLZAK- znwfA{#=nuo`r@TalUZY6FyhkR60-4@QER%$W@WI)ViD`REO5Vy3Pb{b_aiYBw;jsH zhl9m51C!(AjVtiOblT(a{&+1!Cug?fke#JwV_&V;yoXx!e>T*m+;-Gt_xTb9wj?I!Edm8()f@d}M@n(#SEd2qefOiJ``F6k-03&?J)NMutNQ{-3 z%9AX9TBv;zDPCmckti$L@uBBqzuhmT87r!7{n&d{zolg#X773!I#R%UO2#%QIjuwR z$_ol{5)p2@W_2E-%HZEtdjrRC3t)w>)6phOUjKFdLR?_O6+`0MD~@>Rv&h^KUpcTp z0`dOc<`Zhou9(iF#vK}`##(0sn1S+`WfN_g?-EUzd0xon6=l7)zjSsl_y28Y7tY6g z%gP!P9&)_HZ)_~@6_LBSWHOGWQE0kTc#iZ9m zq?v+jg|;)GM^G?2^%l0AGttZ@jJ$m5lir~*PL=APZL&OJmyc!P%+k`~PuZ#U;gpv5 z5&eS{<~n@v$6R9N&w&T7cnOveTh|`6hQZSFFDIcEuFhyX8;4v60xOo~S4!$Z)l>c( zi6w@p%#-yiop{#u=2S4B6t@edeW>+%F6crvOa^r5%Uxt>8{D2>Kn5S!5!KurGx5-! zBQsL4;L)jMCwI5x6tH(S>RVynpX}WU;odwnn5Exbno|Cgy}K3_=vvDnaE@WHjoQpK ze!*uGTtgYC7}O^f(?@YF`vxur@6Of_ZgNP-?cTRj^B0(!xEZtXR%*BOjt8a!L&yrH zgM||#Z~?b;7KaC;1*k2UEAheC9wnvBN45*lM=G@XO+N^>#&)j+WMr7}jR!rR_sawx zFYNT>#9&8-%9?uL37?Mf5(DcBk?P^H0uhA(4{z8j4_Y{}-GVaYFnc;%oyXsY?;Q0I zFvZP-z^%!JlL6hod!L$WyH>&jB}9VIuOicEJFpCv#3>5o>g%pQy7*f^>ITtpS~hy= zfqN|YV00y0L+|p``xYSvo%Sn5YjXtTQ>*rC$~OOnxN_#z?j>QV%_9iHd8#4!f=K5? z!utM?0_%Z6&1&YpU^o8xTjGxb_=cm1Q;%SMN&d9l;i7Kf+z(Nik2f4-)DncIdA6Tlq!C&i^W#bYi`pW3x7VhV>6j z;Y+v~t)^67e1v}Jp%(`uhHyqaC%g>wriDfl%0l`)e0TK=S|mNz$x3+hJm}h1QVsTUi_o|f>PRVH zBVs{EO2Q3KW2GBL`(~$)N!^b+Xv#pv8^3sZC%8~m#699b-h?R`qut)&u1^*HisDj|JbC5&UGrZj{v6(+}H~#elw)MJ16u}UAv?MR^hKQ{(D9B2k(~~17|HjG4MzH6YO_a)tKX(>m8r7;O;sR zSd@8J=9a6v3ULFz){Bz z|Afz~a)hQAMRh7#PT!Zp&lfBquOb$%OjTs?f@a`W+7(;wT4t$L98fC0c|(wYZ*_~^ zM>oT}`Din}Y?CmBs(jqxHj|VXC?odz7?7;o1pI4zJ(P0v0O_f&ShygcMM#ZRnW%AB zMPT@GzQ*sHp;_$=grCz{TL5#VE$ils)aGZeH&R4yD{>aWmD$j1{}dJf>mf0K*xGdH zZ{At=Z0soZ(7A?VIxfOilNikOZ$;w%!d$2=*as?x9nW0J>1$fol-Tgo!KoNJ)an1uZI7GDpL-<~>eWQRA_<{Z$3%2X0J2^@@!Q@4E zF(AAaBKWNnJi#PS2Qo=*Bk?GC{vw{B2MH1SwkF$vcFCW&F98WChTJT2sp-Hx`ig-_ z4_OaBnEF=MJ!W{!Wsh}iIIX0e)$3LhXu`gsmNxzB?2#gv(}DWkBumi!|ySHo5SPM?|kJz2;P|Gh2%YP7`~~ZdD=Ix@`ZekGFo<0)z4m_ zDNat~{HGoW~##9eb9;y z->V}tsmN0K-CFLQ3NfUXtw#mE7=jZ#rc>7L+gr#kUG3k>TAL^gNpEC@0JWzJ9_+zf z6$~U7)2cSNLO(*!XmTZ5!ujGS$%R11{~EA~_ui6RU6fCf>)!8@Yx*Q4xzfLoTw+3? zI>}4R^Qe#CkQF(4ux-(WS~f?=drTv4L(aD26Bi)@DEt8Z2NDk4hTn9K>qVpkjC8s6 zx=wnBWsa|5*n1d>e*G8GYMFxxaTSD1JwfK{yb(bL#B`;3DQ}!92__y<3ag25x$)jr zyV+NT*LJiiY6PVfJp&BymMVncIj+?Wz)n~)!1;)OGKM|lA-)B;BD|1c+SX4?c$~O$ zh7}vAW}7I4kJK$N&6s_Lm2Y`h?5dt$&7r;4`idcu+aUT5P6*RxFKHcr2c+(FDj~qO zT_|1zKQ(42(izb^DK%qs5i3E?awoNg1BkVRdZ`}X}3suwk8 zrQwem7$Nw|8LWmG8+-E&X*oz;!Qn!ficiHJ)yNg6SBgq_edlJ@0TAH()3Z$Q`mV?1 zBaDUp@>2!E0%m~uxGdm9w>@-MtaeB+Z>Wz3m)Z)*&b5b zQz90(u~_>O?S~)m#*WOAXO6*su*h4#`HVS~9Pv`cb(|k|g1Vf)q^3_FK$f)JO&T5w ztZKAdx|qsDZI*qcUJK2^TzYA?*7XS72M}fDLd6z(BGlh(aHX9i?t*q&oByVxs-$4A zb%}V{48&sDSz>M11ta`JBA)-GI zB-iqX6vMRZ~MxS+A8`D^;TslXd5t&haR3ktE8#dPLVJy>NF=Xn9a#Mzl3eib5 zd{?5vs%a$o`X3{@)B@q~zli9*P)Q|VV;g}M$z^`PFa2N^?ic*-S&b)0Q5YlvX2H({ z>S5TgW{HVgi|5Mbb_4_x%azp2l+-U|oHc~_phmG_oV0?BBqwi;!MTg@ZeLSAR6~P> zNE0Xx_{EayjWc6Q5&Nh|YlqA{VUNe(Yd_l?9~v1vxNRwaSef&h{q7gUkq9KDSsM5@N%QD5Gd5j z5Rvt7Ib`g1ysD~}%y$L1t=^RdDj^#J@{E^#6#S<6Cnvo0BT2Ei<5~fJi8xz>b*NVhg8Ki3sbU75Yvys$u{t6<-@c#N$M@QIuyj9!(H0NL2R zn!fA>6>z;{z};0}=*05M(9*V zuAGjy(TbSSV+V!=Fcn8-WfprY->%L&2KG7TNr0sA`e-n_>NZ!)bFDEMAvditKvoT`-1k?3)YBwj;%>uoA zv!3R`B)Y-LHrGP6mul{d3W#b@NTwCrX+^JnEvPXLmh=-VPr%g_l@Sp+OGnfO7@#4f z=zlW1|HA|NcU2z!_`^FnNf*&pdh&irSw^L+KWUEi9UeYRqUocZoXm?UgVZI z`e@MhXE-3M%)q(hgM&XX2J-sHe7VQJG&Fn#xni%Az@F=L=LyIt1F_wK?`!~Gk6THE zG}}_$&i6ibz}Heht|c#xF0@VllA3pm;JcLWlpkEc$gbDuyeoM4yuUcA4-tV1H+f)f zKKEUoT2{Er6OoSI;RMTozG1~^)cp&!1rtqDs_vSAHnZnn4;J_u+}~>iAUt%CCYUaI zzuu1A%D_Pn7pZ}YXG#i4j>u!u(k#{|^4sYhW&hMnckj%49enhft&Y>KniS4X`XQC5 z(WqRKr>uJy)b(|`G(#g77uduvWLdG>Wp{s&oBI0jA=@OZEg4eTGx5;7W#p0fjA4M+ zG&>QOCFYz$w2O4A2#E}=DGYzG(U8H^nJm#ps=S{gVa?)BlK6mW^{+ne4aX-&o zC+fhXY?@Y&g}Z^ijOGd6kkC){-n)*NklA)4Jw-O#cxsMcPwXTscU*T?Q(jZZ*oAz> z{fj_*ybOzEl|G9*HlBVQW6W{T?5s5#EO^!3QGUCB9)T7)=b8Lm6O_IPBXzMOlEQ8X^g{L_&G3XdcoAr9o)-4+9$;U_9y$SE*f8sL|+dcHF=*-c_adOLXU zJ~DSyLrrXGEt^R?!WAepb=)?^a_}E&z5ai9b_dOkxF^o7`F8q%KJUp@M;NZ%Z6Hrz zVY=h{IsMv=Z;N23_U&ra&k2_9d+-ZqezWOy zoODa-Wt*}W9=Uj>mSWwyp|vih?yA!#o%(xJNaLdmqe_*oOgCe{`rVshl_fn-W52Hi z&eFnWPSh(cX-}z;Ns#$`ypuahdS4}t)R8xNCLJl3`7+z2)vqp!PqU1sWo}A{=qB_N+Ii8;y6~@S+qA;imOmkaoml!aJL`V2hUP* zm)s%(ZtMcpM}6-6sg5KC5`Iy|-3x(r3DdRO6srC%Ck&=5!PgaCK^lCzebG5tWA!8Z z9c*RKNf~7JBw|~>jk<6LGG`;SnE1PJshzd-=(XL;xxT462DFKzA*r{|K4l7E&_cPcN9ampc4GwzT#td?{OqagJTOwAaCj z9wQUvk%ny`VPO7Yg=e*A?ugPS$EoH|j+1ejYBsUtQq3!Qq;}%6e&QaS#zl9|)tTTf zk5pFoO4MPIvr`u=&8N>&TbT*BL8?eS3Iutu8Q^{ZoV(+!6UUv72RAY15*m4`A{z-1 z^X@1(lf?TFFi$Wqx;xs+RfRq*=Zm@3N<@?nhQx>~+oYtRW{9@dUDFIx!c@xxo_E63 z1iwgQb%dmldfjKwDM?KJ3Sx60`Ss=^0UT)-7Z?^(1zNXrJ%g$aMY9QyiU^{&RkzL8f~sIpL? zFXPs`LJa5@*8Z|yvnSSgJ>ZE$(CV=BVGxF_;;c8u@E)aKtBH?ST;*VsVC}N_t-5vW=@}AF{{&1aE*KVj6XPi zrb8ub<3O8pnED0WjN-L~07t&Cm@$kg%s0!lBrD`sQ9;wA*ibQwOzpeT$q%y?YUsjh z2K*=covE@3Leqm4Ub`@9^M0rw_t@<~tE4Wy7s{R7Ff(9)b|G8gfAx(@B9z$bD$Ba0 zy_x`M=jE0ttaIit{hQT%qm8eif*cg-whq~A*AmYe7cF~xU9hg%Ts}vJh*8~U{RjK0 zl_QJ8!&5?RG4dM_pik>^Om~n@QcYI_;4dE1zhtdxQZfB5JwP%FPB`!b5hy)qbC zq_&WPgIKKwgp$zs*gdFTSd~nxeO?eAQn~(u)VJN91}rWR zTu?zu&KaeJz4Y=Q^mMT0gOCDTSlCdYXArOP@lcaSw8NDiU#)&) zTlAO4($ZAxJl4bZ@U^mgu5Qt;%r~T|?tWA;CE_zP3H12vF~3qrP~OZFA=&BqK`Sec zOh<*11g{MT0J|93{kba%LMQn?cSyA#L!X)FMDfe*8oIc3@+H@t;uVzJz?Y2sn>bpW z|CnCWmvDr+_S6Wb1bwFvbYK>^vbM9OBTooxy;m-seI<_TERs^s$~EA>uW^vlm$l(3 ztLdZy^#{0K3~hxRD?QFrnt9 #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1000000;// 1 << 8; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int* a = new int[SIZE]; int* b = new int[SIZE]; diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 0d06ff8..0b4719d 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -19,7 +19,7 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { // Compute exclusive prefix sum - // timer().startCpuTimer(); + //timer().startCpuTimer(); // TODO if (n < 0) { return; @@ -28,7 +28,7 @@ namespace StreamCompaction { for (int i = 1; i < n; i++) { odata[i] = odata[i - 1] + idata[i - 1]; } - // timer().endCpuTimer(); + //timer().endCpuTimer(); } /** diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 98fb605..fef098a 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -53,30 +53,44 @@ namespace StreamCompaction { data[n - 1] = 0; } + __global__ void formatFinalData(int n, int* odata, const int* idata) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + odata[index] = idata[index]; + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - //timer().startGpuTimer(); - // TODO + // Calculate the number of elements the input can be treated as an array with a power of two elements + int kernelInvokeCount = ilog2ceil(n); + int n2 = (int)pow(2, kernelInvokeCount); + + // Declare, allocate, and transfer data on gpu from cpu + int* dev_odata; + int* dev_odata2; + + cudaMalloc((void**)&dev_odata, n2 * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed!"); + + cudaMalloc((void**)&dev_odata2, n2 * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata2 failed!"); + + cudaMemcpy(dev_odata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy idata to dev_odata failed!"); + if (n < 1) { return; } - // Calculate the number of elements the input can be treated as an array with a power of two elements - int kernelInvokeCount = ilog2ceil(n); - int n2 = (int) pow(2, kernelInvokeCount); + timer().startGpuTimer(); int blockSize = 128; dim3 blockCount((n2 + blockSize - 1) / blockSize); - // Declare, allocate, and transfer data on gpu from cpu - int* dev_odata; - cudaMalloc((void**)&dev_odata, n2 * sizeof(int)); - checkCUDAError("cudaMalloc dev_odata failed!"); - cudaMemcpy(dev_odata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); - checkCUDAError("cudaMemcpy dev_odata failed!"); - // Format input data (pad 0s to the closest power of two elements, inclusively) StreamCompaction::Common::formatInitData << > > (n, n2, dev_odata); @@ -90,19 +104,21 @@ namespace StreamCompaction { for (int i = kernelInvokeCount - 1; i >= 0; i--) { int offset = (int) pow(2, i + 1); downSweep << > > (n2, offset, dev_odata); - // Transfer data from gpu to cpu - cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); - checkCUDAError("cudaMemcpy dev_odata failed!"); } + formatFinalData << < blockCount, blockSize >> > (n, dev_odata2, dev_odata); + + timer().endGpuTimer(); + // Transfer data from gpu to cpu - cudaMemcpy(odata, dev_odata, sizeof(int) * n, cudaMemcpyDeviceToHost); - checkCUDAError("cudaMemcpy dev_odata failed!"); + cudaMemcpy(odata, dev_odata2, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy dev_odata to odata failed!"); cudaFree(dev_odata); checkCUDAError("cudaFree dev_odata failed!"); - //timer().endGpuTimer(); + cudaFree(dev_odata2); + checkCUDAError("cudaFree dev_odata2 failed!"); } /** @@ -115,7 +131,7 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - timer().startGpuTimer(); + //timer().startGpuTimer(); // TODO @@ -173,7 +189,7 @@ namespace StreamCompaction { int remaining = bools[n - 1] == 1? indices[n - 1] : indices[n - 1] - 1; remaining++; - timer().endGpuTimer(); + //timer().endGpuTimer(); return remaining; } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index c917626..d7b4e39 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -44,17 +44,12 @@ namespace StreamCompaction { // Careful with non-power of 2 void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO if (n < 1) { return; } // Calculate the number of elements the input can be treated as an array with a power of two elements int kernelInvokeCount = ilog2ceil(n); int n2 = pow(2, kernelInvokeCount); - - int blockSize = 128; - dim3 blockCount((n2 + blockSize - 1) / blockSize); // Declare data to be on the gpu int* dev_odata; @@ -72,6 +67,11 @@ namespace StreamCompaction { cudaMemcpy(dev_idata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); checkCUDAError("cudaMemcpy dev_idata failed!"); + timer().startGpuTimer(); + + int blockSize = 128; + dim3 blockCount((n2 + blockSize - 1) / blockSize); + // Format input data (pad 0s to the closest power of two elements, inclusively) StreamCompaction::Common::formatInitData << > > (n, n2, dev_idata); @@ -96,7 +96,7 @@ namespace StreamCompaction { cudaFree(dev_odata); checkCUDAError("cudaFree dev_odata failed!"); cudaFree(dev_idata); - checkCUDAError("cudaFree dev_tdata failed!"); + checkCUDAError("cudaFree dev_idata failed!"); // Calculate the number of blocks and threads per block timer().endGpuTimer(); diff --git a/stream_compaction/radix.cu b/stream_compaction/radix.cu index 4409c0e..92d5b2e 100644 --- a/stream_compaction/radix.cu +++ b/stream_compaction/radix.cu @@ -105,19 +105,11 @@ namespace StreamCompaction { std::cout << "largest number of bits: " << largestNumBit << std::endl; // Compute array which is true/false for bit n for (int i = 0; i < largestNumBit; i++) { - std::cout << "Bit: " << i << std::endl; computeBoolsForBits << > > (n, i, dev_idata, dev_bools0, dev_bools1); StreamCompaction::Efficient::scan(n, dev_falseAddresses, dev_bools0); computeTotalFalses << > > (n, dev_totalFalses, dev_bools0, dev_falseAddresses); computeAddressForTrueKeys << > > (n, dev_totalFalses, dev_trueAddresses, dev_falseAddresses); computeAddresses << > > (n, dev_bools1, dev_trueAddresses, dev_falseAddresses, dev_addresses); - - /*cudaMemcpy(odata, dev_addresses, sizeof(int) * n, cudaMemcpyDeviceToHost); - checkCUDAError("cudaMemcpy dev_odata failed!"); - - std::cout << "Add: " << std::endl; - for (int j = 0; j < n; j++) std::cout << i << ": " << odata[j] << std::endl;*/ - scatter << > > (n, dev_odata, dev_idata, dev_addresses); int* temp = dev_idata; @@ -131,13 +123,21 @@ namespace StreamCompaction { // Cleanup cudaFree(dev_addresses); + checkCUDAError("cudaFree dev_addresses failed!"); cudaFree(dev_bools0); + checkCUDAError("cudaFree dev_bools0 failed!"); cudaFree(dev_bools1); + checkCUDAError("cudaFree dev_bools1 failed!"); cudaFree(dev_falseAddresses); + checkCUDAError("cudaFree dev_falseAddresses failed!"); cudaFree(dev_idata); + checkCUDAError("cudaFree dev_idata failed!"); cudaFree(dev_odata); + checkCUDAError("cudaFree dev_odata failed!"); cudaFree(dev_trueAddresses); + checkCUDAError("cudaFree dev_trueAddresses failed!"); cudaFree(dev_totalFalses); + checkCUDAError("cudaFree dev_totalFalses failed!"); } } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index df6d4e6..b275387 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -19,18 +19,17 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startGpuTimer(); - // TODO use `thrust::exclusive_scan` - // example: for device_vectors dv_in and dv_out: - // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); - + thrust::host_vector h_in(idata, idata + n); thrust::device_vector dv_in = h_in; thrust::device_vector dv_out(n); thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + timer().endGpuTimer(); + odata = thrust::copy(dv_out.begin(), dv_out.end(), odata); - timer().endGpuTimer(); + } } } From e44227c85a13f948f7e8df2e29567c6823d1aabd Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Tue, 22 Sep 2020 23:37:02 -0400 Subject: [PATCH 06/10] updated pictures --- README.md | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 7ab4d51..58e3c5e 100644 --- a/README.md +++ b/README.md @@ -8,7 +8,7 @@ CUDA Stream Compaction * Tested on: Windows 10, i7-8750H @ 2.20GHz 22GB, GTX 1070 # Scan, Stream Compaction and Radix Sort -![].(images/main.png) +![].(img/main.png) ## Performance Analysis @@ -16,10 +16,11 @@ The runtime of the naive method on GPU is worst, followed by that of CPU, then G Both the naive and work-efficient method on the GPU only rely on global memory instead of ultilizing shared memory, which potentially lead to much slower runtime than thrust method (and CPU for the naive method.) Naive and work-efficient methods also do not use up the potential of warp partitioning or memory coalescing, which are features that thrust may have to reduce its runtime. -![](images/Scan runtime.png) +![](img/Scan runtime.png) The same explanation above is applicable to the below as well. -![](images/Stream compaction runtime.png) + +![](img/Stream compaction runtime.png) Commparisons of GPU Scan implementations: \ No newline at end of file From 58fb14c95f4b53e0d52a6935232effd385b3aef2 Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Tue, 22 Sep 2020 23:38:10 -0400 Subject: [PATCH 07/10] edited images --- README.md | 6 +++--- img/Scan runtime.png | Bin 22268 -> 21442 bytes 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/README.md b/README.md index 58e3c5e..6922f9c 100644 --- a/README.md +++ b/README.md @@ -8,7 +8,7 @@ CUDA Stream Compaction * Tested on: Windows 10, i7-8750H @ 2.20GHz 22GB, GTX 1070 # Scan, Stream Compaction and Radix Sort -![].(img/main.png) +![].(img/main.PNG) ## Performance Analysis @@ -16,11 +16,11 @@ The runtime of the naive method on GPU is worst, followed by that of CPU, then G Both the naive and work-efficient method on the GPU only rely on global memory instead of ultilizing shared memory, which potentially lead to much slower runtime than thrust method (and CPU for the naive method.) Naive and work-efficient methods also do not use up the potential of warp partitioning or memory coalescing, which are features that thrust may have to reduce its runtime. -![](img/Scan runtime.png) +![](img/Scan runtime.PNG) The same explanation above is applicable to the below as well. -![](img/Stream compaction runtime.png) +![](img/Stream compaction runtime.PNG) Commparisons of GPU Scan implementations: \ No newline at end of file diff --git a/img/Scan runtime.png b/img/Scan runtime.png index 143eb9d9c4cddd1ca193d7dbae932f387403d94a..bcfd3121f12912bdcc6c012275035e004bd1b7a4 100644 GIT binary patch literal 21442 zcmbq*c{mj8_jg1>DMV#0*~gN~9!epUWo%;^i3eG-GxnvZ5Gq?)3lqZ_V_%0!ma=Bg zHnQ(z9lQ4)J~GEe(~^G%Pg7jvYI#dhagu z*fDb0v14SeCr^N9#wHO($ByZQs@}b=>p8HPOzj)9QG8HL+^I1t&LNmWY>nqTZRdML zpvC$q{p^quL0hG>CS_sF;ZK!r3!YPArl+vfxj{p5{E_+})Q=w4R)^@}k8v4G6f8>z zq;Mv?@m_!bK3FJvVt3z{aOMxf`)@5io+Be)T-$uwE|R7{R@`)vjhyrc1vBXkDe?ET}toz^3U&4}}`G0x-mKhm%=Re;ch$l__*Yo80uG1R3zrLKweOTk) zn`0pAv(aPlWT?!!#ADG20)dQwrMi6O$`?9znTfR@Dx7|Q$oXs*EiBm6odsjKo`bng zwK$u9$bPc$BVQj1t?f92pH2;M5UtrO$TciHLE2FuUM3{zJarHO=VJ_ozD<^Lu2oZ4 z7v|$@EO|0$^^EF_$L@Ji(-!c_=cieRR!~Buy?wYyiw>bZFB60{uMZ2{jp9^T+Kgfm z4@-|Ct^Sr-$Q^U@G{%$jeR+Ak^Q&uX!XhGVC+OH7i{jSb;tqD^<6gi1X^>YOaOfW| zvOBGCIBQ$EvpzdsH!#(j=$A5$I(+*fJS?oB+Hbdl*yF0l$xO${d*3N_e5NB^;*YOq zr<)$UzgEAn-Z&&yu;ybQd3irYPmY`&40GrqA{}EFIFemyPfR#iLwy|pG|v% z@lV8v(9ry`F>_)PQaygaE)QpwSG5|J6R=hlFG)#B$;`>Q99bEeNpJ-#iAzju?guAx zR)Bo|*XWncv2anN!-M@N;|&q|y1M69#f@4sD-E_Ki+wvU$xn#sgU~yPcvQEeb?+kl9bp_F=MP* z(8#z!SUTcC z7kjrq=s1<&-M;e=eMeIyeYc$Wbl*ltC$DkkzfMUpUD=tLdCaMu-pFmsy}s@{oJ>5} zU9N1L>Q(kPvtBcPE7}U%3d+r ziWf38pPZar@?sEDdm_+duuUf|4fSr_+g^KiAR;a7okF2*sjXv}mGS9QMI+DaS|*rr ziT%o;0b>wTn>8vD-`Sp{pS$Ws;7j!#~(AbeLS8jE7 zcKVt^Zf7($zPiX*g736OpFWxsgU>cG5JxPT5X(%29eZ<)r#bpk6;VTGxi@YmCfx2K1<|YG4!+6oK`py*dk;= zBG?7Xt$cG<8CNqG+)bnC4p)g}Lm`4T2tLas^?vMvn$t1U^KW;ej8plMeUAeS#04{> z?pgP&)6vU{u1ran5e_}oLynTjDS&-#YI4ztzZ)1OFsFB>I};4 zVH)-mdGi(KjRS%l-qzTRTGRL?{T1c1n2(W_)j8f_UOTq4A|z2iHVfj>#dDxBDi0K6 zS=K+uaA|(@o@~9M%Pql8@f4Tm?bs8Khc0K|k}L1txwVh76~tEArE^1{ZG4u22O-Bz zw?#CU4@`QVFq7=aogW(v4R*Q5ofJbJ1JhCvC^k9C4J-PQ!x>W-8XRuQm;V!fSH$_F zxu={8CU)l!)b)xv>`fb$D5o5|Af9ezMUlkn)Vt}`iG)Kpinz|y{QP_o%k0*!S)2$^ z`}xVZSx<;1|C?-&^PizrV7~9S{{QjfFaHmQ`x{5H&d|`%2wlG(eH?hEKs-&<+qVqr zYH9_l(Y*azE8~mfQ3@r%wHf*Q`zQMb4Go#F%6gzJt*v+Wva-JV zP3@dD+1ReEz1?_8P_PN;licoVvncC=t9^e#y;aM*&plvTE(0@RqoS|E`m?jM`bI{! zd>TnuO@;k#a|?^ysw%c1_qUqz+Zl>x;?`DeNjAV>4DV8w zuLDP4`Y^7&TbtJU+xVZ~64sVSBZh}0 z1$p^6X!5qr1>qEzgo#&Q_g^0$H$=+`$X|Z(2%ssG!u+y#i}8hf8Q}H3G24v|&Ab8>gnx>Oi9Mo4J4HRnTE1}K z%u8!dTl_u)KI{$K*KCTogHzlNEWtOKnBKfk9JkvZ*6m%?ucX1tebG zS-++*J5om~wbj_Hhuic>H(lo4s4~Nbm=vf*Na6EcG=5d6!RUL;r#EMzxAODS>?+qldt02H}@ao_9qxU6+rA1wJz zmJ&RdLPLclpO0*4_xWDE8^L;29*YBdmiOfr&#wE{D~>hk+%xTQrrGEh98rP|d-kjo z=e8Nl!ctZExkx9HM#+3bj+|NnxIN%`kKA_<MJKm?a`#~9SUUb|qCe%|EHlfxwK znN5Ru5CK+>N6IcxDQpD|gMd}SyzXUKxS_Xqg&(nUB=?6=(c?!ody}H1jKU#6Uq8;r zqtuO;dH|w& zWno%pX>HKdRJ>*JO#BN@|FR~F(#IWKyBm1{aO=Yb+ubsds^O!?`zJSl4tkN*$6EXN znp88z1|Oz4;x^)jU7CgCVq*RP(LTs+Br5ke2NxqG3BL>{ir#sNczpZz%>!gtg9Vn& zb4q5h>ej~^nA!lh?OxNo zySy>kIkVRD{IyhEycME-iGeg7)wD-xTIZN;4e_+(kQpfwrjK_Qz}l{nws+~$Fds7o zlR?Up`{tgC<8`5o9{rZ_{cXvzejM@7J{*3gVt{dQc+Xreeckf{-{gCM`Y_g1I2p-K zze%|Za*})=Uf!L>u4zKPg-UJ!IJpPOGCihVld*!7y|T+y;}m&gF7vsiGNZ=*m%YY9 z`Rr?hspS>x*EkqBa*c{zxF)~xrMih>POC}oNo19++FPdLF64`yY-Nd|-Pn=#( zSBctm&XnI@*Qwl|N?ZhyMX3shSKYO1*Y?=Z*<&B3_Ya6Qh$x#>iqh%5%IOpzL4WJ3 z^cUz3OS2rcoG1IAJ4R+R1tFIYI@Lv`S0MSw#yXICE-Y0pmniIwo>_8lju%B72CZSO zT3HaMZXWWc_$)~yOpHqGXOq?k6P605`njqf1RC9O`M(pKM@Y`{l-kDhwSmcdIYo^L?wp?#d z@;KF0DEWd1GgR_xA=*}mkIUk0ZesKH+}V(mJg9 zYha~jymn48d>!D^?##DnG;+7~u>KUxgDD!``gT0d0IHQsJj0D$FCXzE!ime9Jq8cm zxqe*17HKNVh>Q17XJ=FNzPfmQA&82@snHT-t(!H6yEP?teIjf^x6E!2GvKnN?1PZP z+h#UjiAmQi!*)j=Ns%`f!d<@cU^H^7ZLzCq>9yx!e5GR(4=#^f2dR6Ho35#+++ee4 zjmiAT;F`ZVh^d&!QQ7A{Lao%|zfl5GV0CqM=x}v1>9JVn$OgOo*3aWU>m4c(!kzG8 zWW4EI@V8m5k^rYI^7#qY_6K8wX3d`+Qh^mJ3a`dc25; zc-9V2`?AmeM&EIomoYIhU5v0)VqOYfhK=csSk`GOyRWY=x4og^#Xhb7=z@DoY%k{F zK3bN(Nz+&tGmJ2z>YkKhaXDyGI0M9bAJ3IL{86?w7M^@Pr(2Mm_AE>*lbbhQSKUmQ zz+q(eZT~wQ8f%Aj``kU2Mk&zE#m9Z(Cc@wAJNd=Om$OAeLS1usur%lSg8kd0&)b&E ztH%kUR>n++m*IwHvwp=`{@gKo7%9FUcJG6WoaSK`f@vY~VNp?tJRRXE=yUVEtrezX z*J9P}H>y31FcpQ+uNS(XR?Jalpqllw|N9$H1Ekg0+6wKWMRGrxd&)m60t2|DVsDt`$yFQG2WFRwSAKm!H z85?Y?HBfA4O?8IDZT2;Y=qqo$%)GVh#xM5{n#+PNU>q%`(el={iFkkATIM6ln}wQ9 zct`Dc&n6KgguR)1TYx?a57o}I#jJa6E)Gt!p%Wb{1U}%Q==FhlOl2GLRETZq`nLOE z6~Ur2Lv^z-b-zVeE0YK7g3!x5ks~Dzw2zg19%!>$YP;+CaX|3EPsb8npVnV~hi9-= za-RAAv8%#WjAve}*cN+@Hm7fF0wjKsnVm?8J{k++;>s~Z6uvV+LdJb=!UR%px+u%5 znPOtsNC_xEz3`7^wIH`QcKvdgniBRZ7TN*Vb}4jw)E$Cx;Y6E-Coy9HVbdWT=Y0mQ z6?>2^TtG#nRX3YF#R0P{r9FDD}d;hnKMvTC=jUU)=YJ!TE!WPkte<|^l>Ms;XxWpj zjJ8|}gz6_+&avC})7Yg;!7ORK=@5}K1S-4X) zPjcm6>hc!iJa<(%OL42yrR^e!z_vd2KNIwgGKAZ+Le=bUaAk^(S<&;h74N+Z)h}{> z5}`bSpB=TL=Xn2zCDn67uboh9#Rby|xiw$$rOj|73sv8O#08 z=l{hD|39(C`nYS}>yPa^(n6Abmdb%8Swr~wvwoER@r|K!XD(+k$+ayXaJ#Taj~>Bx zNnAieF5aldA4l@lB{p5W>gwu+gN`-BCr(}Xxd{lj{*d#gk3kaUx(klm!f$7ty8I0I z)CWUabki#~=}rJv{zPI~nUr0Bfn|L)FyR5O$wy#E(;@of%kfN+^oOk6xhB$TUEaN5KI*M9e>xr|n2ZIkOhR(unzl@A5p%p2^ zmun4{IwGPrmPbY^JQnBHs;;-+?1Jj(JX^|{%g~Z}8*%i=U3ud3a6TL_!_IsEo!Q9# z5+b*#sQWn``&<|SQmMPWy?vW_dm=rW@+I|IMsxDp(kRL>{X8t%A1742@!%x7>vsJ- zgPNd8s+>;=DY9eQ1&TkmBysNZCfRub(&}ESjwAnAc4%got?e8rJ}7;?16Nna?OIXj zm(7UWDmJ>Op-}?3^7;P4OnAWlhC}KcDn{Z}Xeg>o2McJ6%f$bTleTd8(`2Kvv2y&T z{FnevN47aUvY=qR5$(C=O`-2T)0dYDck58{IM|*Zlp_vVfDAnEN2%|&OC*VXBUI-`8OLr9MeCnEZ#OC3PtVCn!R+*g94tNX-3_V9Z7=<+)^Uv5@&fo5OU_C& z4Dbkne+*=Ep&RZ8e+#mW`7V*$v-pM2`{>8x1In!b7^OQ8BlFI@d|X8u(Zol&wWouT zj2%KZaQyQv(wNGfmJ7|E73q$MjYy&qp^h4}p{Nr^O-@^Ok4)?`@2&fD&pV+Ilkz(8 z{dHBcgZxZ^BiG&r$8a{|0U@q__D*SkyKtmP7jfUxEQf$iQ$KW{RsV zII+ITiQStz1|d|XyaTx>OzOx+8R+IzV`V*pk;Qh4gK*sRb^UizlIE2;S4i(5&GItq zT^U((en(sf66Z*RJ+h*(mt@UAb(!f%7Rx2i3y!cL1T&lY5V7&*4%$=Y8Wrb%cHs2d z>22UDD|V=|lF=mN7+Ag#+mXssAMb<9$uuGNfJ_lK%g|>JK32*A`w2M?>>)@dDpZ|} zx6QqN7AT-cE}Y)Sf))D;Pi{7P3~1=U<(1z0HKv$FMV&>`asTb`dQ>W++*GFOwi3R(gT%%75~O)KiAF%WS^>}_WpJNz-%9XL7hfb9yPD`=@m?sD`)imvoB-P$sS zkdwIR8F1?8u!Y!2_8ol3&^uwW@qQ7=YgZRB{;SqvOEN>w{97zPjh%cDO zREx|Ooe*LKcJ+Hto%Y|+3axu<^O%F7Xa{ApC+-G(K=pqMBtumvCHe`abH4xaxURv{ zV&IkBubI97t(_$0|McR7P@d@{!{gq&XJGn>jn?^4>S^E@q%H>nS4K_6lrbI2{>Wz0 zXtzvH&`k7xlCk4(Pbhf9ZymDf{vVoa`{zX&#V0{*IU~k=bNjE|o_~)~v8|;KaVhuz61IR<0EL7$R?PjEy62vT z@$v@Zv&jJUT$qEYXbQ~+MT>a?u5#F|>w&{VQ_~c<=V*w;lYtwgI0yWOrONv$E*(EY zDb?Yh6*=v%@?T}j?syn1d*)3(Ncw(hxJO!3O7kr2krT)WW#`d_=&Gurv9U3y*%<*yVUrpX*bG2sw1}<-K8Rjav`u!jlgjRk`R-mvtoZO;Ifp{&%m5|BT(Gk zEe|+6=q%)|`bF2m|dvEp()-eEn0X z)UQ~vH=UnW&bI@gc2hwh%(&@*!jnB^B{)lD7Z=@RLjF{s;Zn zjK$+$RshaT9qJY7mk~9ZFCO^}*Vj^-AH=B9Y0-S{ii%H3%@VzZ)|#w09={egDfj-ng;=*t42&fKp&mT{Z;NueJv0$(tw z*@-134Gsa$&SD)DB*$S(|0iTt=>)#u*e+B*-3j(U!IF=+9j+*5)Iy<9WW;&razg~c zi|_YGG#I6|1=SG4Jvg1fMWaft&9}HplB>Bv8RQNsb=Z{n_umdPj@(YdzajN00^tt) zJo$rLW?qGWdnIO`5Ie&uKSZiM+G3AB!2aI4WV;P9G!iQT z=<0x3TWpH=oIu)1(s2cc0(G;Fdx&rcNL)LEtvK?Yj}^()m6iZeRS0aS|GjmJ4T;~~ z+j8C==&M#f7ew-nW?7_A2t=JDaQN{!&7~{6K!*1zd$*@k9sB$HB|vUgxlqtNWC7^j z%?L-|XZ++Oj~wLI^Z8seyFya{esxXL@blgl%NgLxz-NOg^cy4D7b9htu;C&VUr(Ij zSm3bh`V?U|cnq%;_6{!qkoCP-+=`pZ$kH79a&c8Cqb(A^Cux`5*zRdt(7?ic{=C7; zkJp#)pg(*FJeshxN(q4?4gFj0Hhb<} zPCfU*zju5?3?Qdd$}L6)*mN=? z&&&C|n(DL;eiC4Hvhfq%*h_z*sxT4+fTd4zAx)LPH~IB)P1cZ=C`PUonn|v{q!-RC zyb#Klf<15Q`Puq7$ubgX9Pq1k)JqtJ)N?<(*oTLM-dqKy8N5b8Ug+}Wgt6Ef{U9Te z5u&=+-f$V~P`zzcIr}LQZIEZ$TZwS|b4M+}HLQ(mo`zwD;C_4j5Y>Kl5CGr3s>EZp zW<@NmrL>lPG}#ZTe=@rZg%@W|0WDIEP6Mf2GsjRnIMhIbevt_3Gs&cq__XBWS~wT3 zpjl$^B&Kp9x0JvYMmo(Kwm0TAuOEUeL4lZ&NDu2;nyUlrWO;t2v8Jdhc#S0wU$Iy+ zICQGkF<^JlV=~rYK+gcKmw$9jZ|-|(lrYuR?*PDY9O+$6YX}hsLI+9J{aztXPAwC0 zmY`yX;R-=X63TOreGi;;R}~qa1}g5P5YNF4VhR??Y@E@yO=*JwcB(yu#fQk7$u|)XgqOGy@L6by#Rk<^b z&@hi@9~@CG9c*Z5D8CE3W$XckCgKOUNl^wr(gCqhNSaAY3|y-yB0eDVu(kH;(B1%! zgQPBtPF#6I63e2+UO9+DnBV@F6YOUPCtIn))TrQu~i=AGSI-j4N4(8O}z z_y;NCtycH_>fTOVv}w6BWu_jmDXe;6wEe z2uxH|%41SB5gQ_laNDa46*;acPAT}BE}z2QKVa)slM-HYCli_22V&LWZqL|N=j`97 zMqpRhL#*miJqN?PeK-&hlZ34}`a4t4bEqaIE|XWA zR??(fRI}tUwO5HREnaVx%iU_$d|Ck&zr5`EcG0-=b0?!n;ZShiQ6ly-^aUXH`#pDm{(l&)(LfZSN|%pXDBu- z=3RF<9y;6us&^Xjli}ggIV3kC{X{K_&EvE}&sJAv3%maBIc@&fjA08yrQ5-n!Y00VIh7(C?JyWkf9lF!>c z+)V&&lVjbYJMo}%?cx5M!S6vS`t+!Dgb;iDhh)v|3YL^>cf zr?5?J&5U(A%cyV!CYVT4oYv?g>^MeYfS?3%N`^NWy@8v#|HHe*&4yjEN99nja3ff6 zLWzD=&PoOJGpW`%9FKEQNOgHw-^)+Y3yNbVd_)jnO9gFwy%JTXu>|@?M_b65>C9ci zIy-fSdv-TgACTyh%IjK)ljJ}))KSN6GCO^=o5L12W^E)c*`jb8rjHG>`B!c>8~Q(E za(E(x(dCC0(Dg6ov-{QQrPIU4zJI(U;WBY=z9%RCG{FAN<&kznN|37?-!at~cw#eWa|G8AN`QlDG* z*I2RWrMXIk6JAR5W$5-NqG^P1Q3_*;VJbp_x;a)5HdIR^Ii?oUJd@iIVQhvvPI(?D#h>RKO7VCAT6M-mSrTRQFmW#`G02)4g@(VX=}Jv446+A-gkH(4@0AOmb9B@5Pt$ zM*E&O514VXgcFXQ?G++R)`(W$DmAg5$CPnH5O;?FP!3)I4D+vkz~>~K08GuEr)_1C8TWIwzXH1$kj;u}>9Q8!f$a!qHT zLp!0WYx3CUXnGwJcC=#CRCm|ec|wZYL!sFJ!dY<5Zq(@oGj?u>l;fkV$F-CRC#ekiL( zg#B!1M5lkbzIU~W5Y7~%e+je_pjT+nh3)c#*RHEl1NG7Qxc>dNB#9*ZojRl(fqF(h zBR`H9Pv4k5sTPZ-TT-Ks)^Uib$ZC#YDP3gb^PaddER*Y}B{}!Gd^njE=`Gl!_pPJT zM~|w^Z@wOrd-AUv6s2Z55;mVt^4e68oaLHAb#!yNhb4&`jQg$m55 zvRB?AmgWa7*CkT zIcSyd!zehECE=_|<(|IuuH#$Ey@kc0tSCgQi$jcZ&yC5oiBlr7K)6-T$Kx#0_bYNT z5l7*P_#8&#zLe%%_^nx-58<(tK{l%Hu0uD0kt&IGV>I&OrdsF8X=TcU;ps)?VTH^@ z*L2W<;NsBTTQt0}q%3X|JX#B6H4LgG@m=R$ zH9<8nN4M1455`m!)<(=#FKh&cjKmISwb7auklBs6J9NLjW0%aTXJp#s1??lZBUE`U zX+neXwe*IMg&wqhy$lrF2T@vUNo)56e+?WAD&G~Vk zrdVARcVmLIx%=n%919jCskRUX%YDtDW|UP5c9-Fm!l#EMCPYK@jK#Dcl@0MDr?|s( zr9>pvLMA>B@iV36T!K5q&~#5Kqs-V7UgIl3%uAZ^&6UiA*K)v?nEI0KX7O8#F7wr-F6*=oy_E=Y|QM|q5{y9Jk|nlSBzGAbm<=wh%-dN;#uhnTQV zFrDR7BfeYL>0zfISO)df1O&_v6z8A4^U{41Qob90x!1 z`58mls)-tdYeo+TuSuqi*XctBJM|KUybtI|FygJmUZA3a>|dsuq;Nw*pO!%>X+^QG zkToS=nmLR8a!&l%6z?c`w4Y$*d@4Q{{zwdBRn8mdH+80%Z_Koyq${=Hp@xkC@?{0j zl&Hr@i}fOEOMbJR(s6JL8Erq6EnVV0%yIzlT>-4pCp?RcTbBFMP=7nJslEM7#Th3o zS7Wwdaax3(u#ee801kYdUGnDS;t{IYbqA8ReL@5cM@K;?=!4Y&NgHAPi7tTvq783*Z`}spF>3dW{3WqUFEm3(F z0Ka@`=$&pXfp$V|@i4|nz)_bKsV%~C72P4iB8Ao#Ww~V%`?53Z!s)BMGZ#+Z=)K1r z5o4I|!E|_)9u^VRM6SM}d)~w;2=uV&4hdxG(|z0|BhR@Ea$W0&G8w_tzfyIRVJ5L4 zEzc2L=6cHX5-t$a-C`^RAv0+DGCE{YGWV5tTKP2LSBg7)NHl{LIa=?MzV2JD?Txb< za!$E~O(?=MfX<~-QZ&~lM>IiXejK@ZFTJ~mU}9V^n$J2EAwZ>HG$4DS3i^}K2!s9#Zcz_c_4 zH$DL}7qI?JA{tr6B+a)uD}*!Xj~$Si?~99oOqy_)tGdKahIiJvlCoTxFDGuhpb;r9Y|$Ckg{6|1rQ!8Rrd!jF}l+k z-h9#zvIA(Qlu`a_4l#2mT?0KG$*{jhJsk{!C;N|yyQCa;;;ba!>vuO46p=*1O6Z`~ z;eZRXF317|#X8tMD!McrVwS3+S&;*oTHvplu4F%#bhpgZrLA$mc`)C)QhkqU`Cds9 zaKOT`2%4)oq&kG&L%>|CeqlV?tC7T_|E>!m$}U$wTde+Waw9M^#_ip~Ole6q_NHg$ z+H+38@Y&mO!~DY=jb4CE59D=u;{x&u4JX7tD5<29!Y6b%S#uq?%gI>6ng{ZeruTqH zLvw7!2Ms$UsDlCv4jCGZfIiFuuDBv8XRv&T-(4Y|73nZ@#PdB)J1?-24u*W#ucFxSzL<-2V3Xba6sTf$2 zt|=L+qve{StFDMtwR%+Intp>73~n{KFMaU-wS(?7rD%U(@iGV;zI`liix3nRojek( zYt_A24w=X0jLN$RgMcMf$h+$huF9N3r(l(=^^{V`hJSZXiB<2u2Rg_EwG$Z{90gqT z<#O#j%*9Cx=v7Iz#Ks7gYU!rGAgK>@`jYApF|pHs$gXhHlWInO#(X8*6(3fF*tV4> zMR1bVu06A$W8YTezQLg6a$nr}+ZbJV{oOHZojHwZWz}=vBHrGk;Z5T%;07P#uI8@g zZb%4c4Z$CkJ%h7)1NznB37Y?jXYL-P*iIdc)%BwQ3U6gR0B12b@Ka)rzB~Dl#T>V&n$S);F0%VA&cHu!iCBl z$rmAaW^u*48c5=a`ee1`VvZQr@3AbKgnPzk2YaAXlBNcNY~21aZ&~wa{P@CRvZL=o zFM9#Wgv4BM4v;YOW80JV21#(31 zcli)7-4iqZC0=d&alLgW0*(MaEQ5yli?3=yo&nPc`wsM)AHY^vcaIF&Ss00lOW<>x_j19(B1QyFcA>h#6qc z@&f~l_niXfE1fx%l{vOs_9+=Cc^xB=2xzPeG*Y;HBTjkT=64mVq5m~G`tV(FG@)zr z=j-9#-CEgn5C50z0n3#sLuMwTF^T4&QF*HR$pJ|BC#jqk9l3Do17nSqh z0WufEBygyM#U)9kA|_#& z4KMRbRccl)3ew}a=axvno3Lo~z3m#aB)h?!(~_96`ybJ+>7>jDtXQ6{C8G0}Iq=@q zL;Sb{n5j{6GKstgqr*V|)GBYYziS`}0sV34#~WdQ{l|%*k_bL-x<>f*6cZHYO57lo zP6MnnX?o$;n7%$XYo!wE#Y5Q^@izU2w_TBI`os-N*YqOI_72w#mMB&v(Ok@03~0rz zx}>v66*^8f^6baahjo%9k!WvPS;(_IP&|pS#H*;L;SJ0wMWIkEwFNs#dt%@-a#6@}+dXqptp{-*m?bT6{~JMXS&`LRAe!8W zg_gR48T|pu41YH~>V5!YKRBu@fnp;l+elKC$x@WxOHPvkX~_v>e_)jU0D8;0#BO#( z$U+GM!90iPB^#kCUWNiV_yX3r?U@kGC#mMT_Ox6(p%MtHn+%@(owX9^y#}aQQl6EV zQ$-SV{fB#e+zV^#sr5Vz9O=%ZfNYB7WH+zxI_;b?#yNdpH(Uvv%qzo~+95UgcsV51 z3RfnSZS+&Drb1lOWkfx{ok08bm20ch<+?lQ+DI=-!CFE=wfMuuw58c?F-EwgI|zuP z6ec*%SmRxED(~s>;Sh($vl86iu{_+b3ftojTMP7ik|BDD=LPlea^=c@cT^$=$w)JE#Zgf~h0fWPJn@TCHV_Kv-5wZG7_x{Sw;Gt5JifSSf;wgKh=&eG3Tiu%H&9 zU*@73QJJH%qm*4Si&Gg~U$VNW)Gc60u-4UI< zA9+6WjPj82UVi?!#7`U*E8+;KWJgl34oGbJ6cHQuDo8@<+ibl#FKd(6s`k$E8-&WJ zA%O}lCEx-V2<1dh#3V%XL$fuwR8f^6dZ6jhdkqwYa;{8UXzCWJSwj5QeimSJQrbx_ zT((vXx+$>}uIt|A(_WSOXxuutK)kte@T2%}uGymFmEKsOc5#$IpKEW;$}qj5a!uC_ z2IDEsx}@`prY?o`1l6;>1PuW{ZwL^o2%z6|jT3wGLn%|PdUhdMv%A-uYzObSshP%%#bA9k?`YI_9qy%Z?l3ROiB2BYVGchVq za|KPy;V~RCGgb%4Qpb5&v+Z)eTyUEO(*Dii2H#z5;}0nu#jiJ7ycpP-VN28s&}{D3={+kfMBB2Aa6Ba&V3?QRK+`9v05y& zgjg9iJ(m=69Dk-Eq!b3>kyJ}JxzVZQ5OadAIimB}7!dP`3*fI?7Qnp=Fn*0NAN}3n za~5`1jcq?i>8b6EJza1&xpuS}DV z^Ng-9rjffr5VWF@Rnqidd4w;}INFY%t{EePe5syl;lkJ6+xle!*Gy1VZ^@^XdDvi& z(o26*n4{){a`nz-%@pRb+8 zm>K|{V&B2c};-ZsydlUfZ=_-J;p^GzxsN`H8=4Yn+-T4(nNya}=CL!d!vAi@V zSs^Bo-IEQPgq}d7*9o#~h4w3RJ)89<#Kmx)s2ilrV_||9)u6ic);2|V_Z(0hQa{~< z@~no!@Ji$F9!ubE89-T0Pu+}3{|gQR1pbDDa?p!N2HGv3@`gcurPE)4C6;66 zACa~?xE5}IP;qXs2120OOsR(G4Kpe_#K^n`itBr`^;kkQXXSv#&BNoe9UC;nLhF$ZiCpqt zIC44#Xbu`<>MiQwUh<^Bq*r6XXM{SVA}cm~F8DQT@xc!X$=^wXjg={gj|++a0ETm4 zLJfa6{XNCoHhlGi22~J~15*`t(Gh+w8OVbJ>$u&9!!6px;6fp5JVJN_liI&EE$=jo zD~kYSgJHhs@dVbItgre!nLpm(R`lOqDcNJh+%anL@nBXBG6N+z|Lc@JERTlJ#qx^$ zcAxu8{0RvAsZ0edkE@6W_yUVbo+N8(YR0+l*tc8SY8oi%P%h-$gtCD9c@T?TuAm{& zNVkT^m7U+1AN4eXrHiMLm~2L^(Xz7y|3s=Rln?o6iCL z*#5r$2%V*kMUI1SsYBx-q1NTo>@RE5^i6=q+8SO0tko%?&@XQS&&K5jvKKXw#&Eea z<+%Q5uHuYrwPc7pAow^-1=sr9+Ui$_77dK7N-iV2dsGB)({Nd&>^^g&dueyD>tYb9 zhE~v_qVAv}6kRa!;)}<>Qjp^7Cx53P{NnSvT^b-b-lR6k_}#w?e2~mXA-UEb?QN#q z8fknFoEmw+5Up;K+bws;SE=WUBWEGXb$k&&5Rdj;Y}GILRkKikv4%zvbJJLNNKfBi zsCW6g&ENb}L&8Up^33Rx_$S*T)QJ@EuL$&aM=Us^+%(w(Qm?D95oj-~yo0^TH%Mn0 zW4NBJgHqT?$S(Wh_&uIf7c8+HrHw0GwKBLe3BO)!aS#A*eBGR_-S)8p0m>Cq5M3Q} z1}^*fD?qn=$Oqh43<||suM~FF7zd~asV62K-X>^EH=AW~`*Nh^#^NJNyJ-8h11M`W z><@NAE0JC|FLJ*wu4Tz;Hr(-~NpSU&6=ZO?0I>6=QHZ)=VQlumjVb^*&o2T;;WTdH zW%GLzq?XDcCg~WC`}s+4?)ZlHf^3MDh)fGY^DqD%tfR6nXp`1xT4mtl6xnA23NRdA~&Ph#Z-bl@E_RRq8 zn!-4p12(M>>;6P}NZTG_r-dn3Gsz6x*<5qfc5ipI9}ds!AE9zYJhlD#!&!o_BHeF$ z_Ka_xZo^YoA?ZG*-*euaXALWhDz>tSm|o%{utbFUDU$C2xxw&$M9D40%wj`})8#E$-B z$6dUlqU|D8zZ4U}w>Th?NJH?Obj1){X)0>;ERy z*NZ57AEqxAIdnM`?|vdo_Q?(V=;&yKvMzgh3u@&pR<3s&!Xpjz+Vn0>hyPY*&_H|% zYg-4J(9k-W8cYfhInRcR)-VT$B^zqA!u|3_#Hp4Ry3IZ%(Q zK*L<>z%8?>;PZSGqz>la1H~Tf?;7`E;C>vWa+a{T_(l;$TuFUufVYu_8L2IRbUor& zP!Ipb-z;?bn03^%*woYjBl3AovD`&#h24v+M-SsGKQgG!!iv8x^cOD189Si1yg>QC zH%*C_?i9HF0u~)+n=4ZGlRTVs100G0VdkH=fk5XU z>2Y-H2Y4x~TF@bDBKwMT1-ds|4>Aavlkn4e;xA>v7b|T*Xp0ty%CFBX`-7Xdn&^4>5~H3MSeob}JZhV{!Os(KSr)67f>_>&qDI>z%i zWX0NkIwU?Q+kAP?Mg(F;;lDU_yPFWc0mDZpaFBKA^KA#&?cO|&0LgZ;&7+=WK+yR03y6bCGp z=&(u*Gu+Z@r@udmv}2Tds&$cby6U*>Ui-s`@hn`b(&^i3FDHGZZm!(Sy9oHBjx1%1 zB^_?E)oYfdhprz9aU|S&zCy8b+i%*qJ`M`QH;372cmSJ9_g_ZQl^rDRko3;?K}!Y= zdUg*3kH%oe)cL3&gps}_d|3lG1{`qT=21v|h?Ez2=7=@NELTa=ACy1|&|GuuH)(|(>C7uAYb3aC- zYztMngYN`6TYoV|LY(XKz0~=<1Rw)2#aeeWd$>T4?!9kc{=Bjc*GQx3sMB;nDoWR) zm`GIBuV`a;xvti4YVLeEp2h&@FHn8Zi1#0wZ^ph2W4j(~fD3;Jxp?L*CD+mzC$VRQYXFL#6?62NLiiz{>DCp6@a)k<)Ck#;mGiiCJB`fTqGa+$793bEyWW;Cg1j;0>YqtNYBSWEVWJ~7OPU#`x+u&(kaucY$`x0aD9L-6AUHS zaPt~X*F`FuE3OEGW>SLX&G9^9l|l(~RBEA`?d3`T5X0Pj%{7l}t{ZH?Y)^k6klSb;7Zm+pq7^J7a0{4l1Ntk54W_J@y7x`z$zb+>PmmirM3%Z2RJ8=T-3q#$tKk$ttDtF$v|!~upwpuhp%Zh z1N%7vI1?QkR1iE4%Zk+1&I+S6MquH+XFwCe>h=Jtj;woT@4*2GhyMs5p(`sVOP?1k zXCfdfkJOuVW6}A8_@K9J{`mn`&bHP1SZiwr(9xi+D5kcpc(FTeG4O4+st3`$cp2)S z*7^MD)XHFN){^)fZrq4}>*sk1pL9k5;?lqm#k zBH>y$ywV7+0WGDp#mr6*)sq(+pI*WV8D>GJ7@kDKG~M7@p~$|L=I$tv9Vk{*%}wFC zlO`nNlY*<@sB-PyRhNxYf8c(h`I8v#q>7&w8n8tO7j!>dUsp3I^m$?O oC0i1?_rJe?DZ8H*INwm!xX{^)rYC*slaX7zxBGe3c!tyd29ia58~^|S literal 22268 zcma&NbzD?y+cqo+BA}=sAX3uO($We7!Vofak8~s5rJ&M1G^j9icY_KdEiJ9W&>`LM zUJLiWpS|z*``+KVQOBwgyDGg z3i#v~+O^`+CH2p;PadnfjjyHRy2b8Rp6^btoLdf$s5cf8R~8n2w_H)vL75lAN%LM7 z25aBKBIWpCs9bXOY2b&iv?2HEqrMVe4Pf_=3@gFH#EpEU$#C;34MudNclz9kpm@rX zylGFGtXTSjJN4HFo|we_-%ZV&3n}}a9E4#1OXjK^zivYRxK{#QZ-S#;x^$Hdx;{c+ zfa^a$_iMrYug@>}V}norH=kd^#0S6p=kG6JkwBmP^*K8D70!*$lf!M3p&Vsyr^RQ7 ziy0Bdy(xmxk&*X?hleTLZ=?vjDsppkQ&1frtrVMljiiBVYJSB4M>3DVw3@1OmcCXq zVcM6*uBN6&;Y~^*>>~gAIys}T%j(6G+$-UDmq<#C`T5O__=5wtj=nz9kTzN=DXB8M z88I?(pN27Na`LWz*N;Vd_1&Af?2tU(QQ{XQB;299`Rms&-F0!#y@D}u?<1#hv5Q0d zqw)Lu`=#%#PS%zy7PGRm$9?Aw!@qVdHi}LrX@7XZ(l>z<3kxfu&;Dq~BJsiVmmQt#T4k-P z$-=HhhiR3t)FaUy`_ohBQ??LNP8XK!G(pENtND7gL;DmwrlE(YRwD&mqP9m(A3uJC zYiU`em0Jul9v;j_pg-Rb=RLI>*)Pj3D3Jg3=~K7;=xpvrtzjD;GY7}~DtJzFc1%_t z9+hd|b8j+HkBZglMjaLwmWC#2(rqR2(;<`Qku7wWO`q_KAATpZ)aNfICY; z<$1y0OplMsN0f;kv^O>;Y^=^XU*W2Sy77301taXk*j%RBvFnUwmd?q{%uMJG2@af`)SXy)_3|aflQ6NY57L2YVg?(%$6H;TJ+n!u zO)12ZqxkVE&$cITvT0Awuo~jQh$vx)Agd~dG>!=bd!8;?(bCe2oN@G2EN-io43Y3v2Fkg(RJrS7M;BvS^`(g_ z9xv0Ccpb7i-_ktVuCke+NUL^Qf^XMMcUVMFIeLpQGk2U6{8DXO&XkR@E-jIP1+zqZ zcRT_oc*KLt=DDLN(KBg&9&M`O^nTQbRf>nP36F|3DN4D&{Q2iAF_i;ix#JtpiSG#t z3KA5nJ-w#m3eUYc_cy>+lN3z2PgP#R3u8D3tqcO_>O)U0DRBkki|8EU-2MHZu778vcMj zvy%DJGt!NrA=PA?SJl^$!MChxm1hi&1BUP{tGKhe^gFbfa%!IajyB$M%Kp6 zO@9=+FEa9`3ioi&^3L>aOj!S@Rh4{1HS(HTX5Lg_HX$bvUPWXJ*!r);@poYVp9Jym zfDe8CSEl)ABKT(p`up>wS7l|~9bH|skH9VWO1v`i^5lNa&YJXpkm7b-*VQUP_*;)x z$jZTB$#eL(NlCj_$I4M9CEtE-6egm(;=6u+e$z9<%EHoZW=rsKIHkFZbmtD%l`EYM z?%Q)ixhm{?f_tlDbgUXhbCzI(LJfq?_mYon>b14ChNYVH;^GD)MS83>G&FD(l~*l7 zeFG?%4Q9S*x#cj+Se1=V%t<|zU>+qCY8rU&#!(15w4$yGc^_3UM=Z@mP>EWW=^FSP zqFP&9zt%b~81F6*&zqg97#N{jgNWu!54RVRt>18Tv;1hlEr^aL3Nd7_%q5U3y#qROiM0KB+_-O;ngTfiJo+LigX4|B(@;;W~ zh3{@h;gz|mAIF=0o7wamG%+zTjoX5W9ag_{#TS?R6oE9GH1odQaP~`>egrDJFKedH zgW~1Om&?S(M<@Di9wLS3$SUixmq3+%BkU!-DyWd*L>@B{j*+^SC5eQX7$a7Oyi0@d z4^1+wWnWb6Z%zZhEx@A@k(PQxg-^ zZDQhDV$Hq2P4FADp_$x1ohs|Zwf&J^L@l9a6BW7LOn*!Z+&}LVi?Hy7@!jUbd2+{; zuCA`BFV-E=4D-%Mo)z7>v|WRxJ9}e032RgJquCq3mR*My_Lp7N?6cdSHLrkN8+KC? zQb#_CrZ(YFn;fZF&%>o(u0L8Zbgq1?)6CA!7Jkd!L#sJfX%!#v`gPqYCCJ9oEpCLG ziHV7(KUI}Y(ZGDT&3Y-yziN}39ryiIM2v~IYxtw^t51DS!<+Ef3?jm+Hx^|POA9Lg zXqY&1*Aaz6jZq1^Z)M){y{n>m z<@qpnqo+la)XFMIssEZJ93H)9c3LlH%6UyS(_-XtTOh|f^^d<45f+DC+As4Om(0pJ zdXHS~FuTqK!6A+}veTu?HUDv(-0xTM@yksnRBOiKsX&esrK+SJ>|0)bc(^sTimmSb zOD%mSzr9a?Rombd%pyxDno}duvG@#^J|;SwRh(X-{q$ zpt^5m7s0fd$>m(jM0`(5It=81og^YkBqRl{d+M)K>NZCA0P%j#0oha4i!60pB`8 zfuisj&h#FA#4wZo?qOpQ7IF_s17Ov=VLUeWVe_%CJ~UNjwo^i5P2Qf3PXhMtjm=nO zF;LoJveMCgjt!hNF##d2Abm#6%5a8$Dh0D^!gWUqq17Lqw8yx#>B3a9T6mjxMBA`} zb7qX1$v`Z*u_aHbVR_rL5T$%f_;aLC=N5po-#oC990S^Tg99q+>PW3`m25fJhsGmN zA2N$*B{rT8MZeqYU;L0Wc7@oU1ug=RficUcopR0OL1!AZTR#(C}ROh!)!r>Fd}B-6-RFOT$i1WNpZ)TQ4XN%f!kti)h0s}9gVq1_RW{&9 zI#*b^xcW`6=5Go)Ek;fxXs_MeIlPNctT7_(0FE?amhVmaTyIuwW+uN+10!p*98U_0 zwiTBC*fuvbKHUIu@TTSkt_QL058D0z3D_4E!2c?c|Fe|1%I25FFsoZ>Eu_idpUi5a z=MZ;TD~jlN8o$-U1t)wc(>YSeFJ5hCDW`&(mw&5DP0l`?joIC26P1_8lAD`rF|;qw z*%_S>v&n5X1J`oq9Y5)|$9RNzER4!2kQUR@{W_|ql@z(FpH@f#u zOYEmoI`nF~&q4)pD4FfvaqzA%HV%%k$l30w0c@tpyRk`3P*BU~{CKANcV`k4(wh^e zw&N>kgN#ixZ7v20Nx1d9kFdksTlZe^^FyV>4kaJj%QyrTr+Tw1YY%z9l;#x`A#td^ zo^^J0k)yv6tCMuS*;GMS!_w@>U)v-b?;UrBa5gROeFu5z@!9V;;=E(x{$Ip6fYi;a zQ`ifl7Q{L2nTrZ0@*U++A3FE9p55~ z|JBJ@RyAXP))F@?i%13a$C{&UEt+Jk-nBTW`E(|k6_$}+%oWg@7){~sr63M|Unfc} z_LPM!&e><9e!UI`fbRS~j$}B3Z~j7C^`_rGaYulk^#t!Jn9nUNbSVf-(c|2pa&m(Z zurZ4+E*mRn!|#{3lzV7Owx-uMly$YLe~qcE(3c?Dbr`9A4&LkdLy+{&W^VhUKMvIc z+6ou-`N~oK7(c&D6IQ(B=H54m-|5vkMTJuecdisS+lIRuA;xx@pk38`w|z3MvJn%E z6m%E+?6OYGnDwSCMW$-kyA)2k&Uia)_Xu>sI2-8~dQ%-c5w*gMM-4!Oee+yX*}q1Q zHnxJ8(n?Mhm1w__p4okV3N5K1OJIH5A2oCO>W0{Za~kDRR0MOq%UX0hwXbfP*Ou&d z9lK6d!ufa2u~MTqf>+dOc6)_45@rW-RhVL)+pf-H`n9}t>53X077JpXq9W5_LlIBI z0s)fV*KPtuDGMuW5A6r7&{Sbp{@Rv2InmXa8t#dV)Xc1kRsio!zL}-Df2oR`z-K$P ztxEtE6C9u-lbyx>`PxeCk>%lh^q5(CWn(p72r;Yr*QLSi9xl8aHw3iGZqZ49O}1+q z`!TDmk^n8L(~Vjf)>O_gF7Y|rPM_G*Ik)wtP)vFFoGKzHD5{}B#C@Z2zuWYjLaF^p zdUpEbw!C!}<$-$F4e6dCieA{ZtnQM_#=t^mW?q5f#)ce7mT6o@uMY)w`empS&l5DU zk?SGh;mFlhYr33C+=4h(jZVitQM2$akiyZv=lj0sZ|^0IO8{f=?GBbrJ}5DkkkPiA zM`Zu2HMwog#~lkkjT`E+&)#iduAHbRM$eVE!1p$0?1-MO%LRb)hudm| zEyf*)PEJlv-t=)fU)`nU8MP9x-C>O_P$*7JOOk#kw`=IXI`sOPw^zf(Dlw!BZplaX0Gi90*)?mTj(1H1viw_$$M?3$$h#HFL2bjY&}t+r^^ z9pB+1Zq3U6CU%Xn8|>5;hXo%oSzjooVySCA(M-ica(f-Th;c3>)vR+{nysGksb}7B zS{om(x-`P!xw|sj7kfogLc)6HFh)tto>_ZfVBjueNqAUMa!coq{<3B0wUGc^B6;b~ zn3gRkdCPs)vGYky35gAof{+$Gg%F#Pma=oj@fmYb5|)yR$5W*KDDzrvJC543qZI|D zDlGph8%P4YnElLZ&V-t!$Bp=GJvEC{h)T~Ar{~|KR7v&iFk#9VAI^oV&AyjB=OCS| zgHuuJGx@7gk!t~?oKYa(4Ai6W;Bi9hPsP~5vJ-|%ds0SZjklY@C7+=8z}_Rb#Xb)YhR9G8$_espS` z&L=YY7=QvtHP7lDP1*~4xtj7dSLcShPmrCBxHR@+Pymv%;Q#r z88SCym6QZb(!YP#R9P?O4-~+Hd1$A4l;mwUP7n=|Xa*x;mDh}jHOHS8h-$D%KtP^b z6S$oVbIDH4lGR4dEP)EQYK1GTE6^jqsj2;gRLI&~cEX|^LA|D6>Dks@6lqo1sy+|&Lh;x8!JEQrL~|N!7^MLBIUhoFB*yKB- z{r>PJaNg3O@|Bb6>$fz6_YRi?Ge2y1aW+@rvh_S$+Sjnp6TrPj#@)&|M7~l3D!q4l z8H^t|757(P?T(zl)7S6qFbzH7tu}Lw_{FkO*d#Xf_&ejXnEsk%(PIR|^ui!zz}j`b zTWcgouVw6q1(sh4%{;yelH_kkhY+i|&n6}uT`$M4+F4}8yx-9+md`Oh{( z3xo5UEM=`_3uN&ok8J|tlUNi<1mATYDGGSf6<^9^q^MU8vfT0qfbqlG#GWPrpR^#*rcTGK`hn*!2qYqh*Dsrk+8*K*m4sn*AP>r<=( z({34mess0KP7n}ylb#!mP`SRn4JfTlrF5~XQ*F$K#SbBwH8tO7#mrBQO-&tczQw&s zWQZxI!fW#~<#cdE*Y+yg4J>Tz4!%)cqq1iZWfu`fC8~Yp+MOT28vrWOEGsV;U3YYa zn+ryVUHs`euNsPwp)I*9@#v8sgzrcI&r8Il$g8kXS*{8yHLCr`D>6zX_ZT0^zr}T% zYaq$_PQ`?s@mkJhwD)Sk>vugXEU^{#-k_>>0RA(vw>DPZK53KVHjbm4>6dVm=pN%4 zOl|o&bDF4{gNM^;b)-i{jIp>6HD;H_4#%TP?WeM`QJ+3ZkZ~KoopfJFHBH}Bu^4;s z@Zoot>>ocqzG_^*%mIhCzDaccK}cLlPV(iPZgU#JFXn2c8_c+PvCHLVncVJ7nE4!` zclqRs>eq#8w>nwT0Bc0fwnt3t*^HG***2fL%tj4^C>HZ`%Y^zx6sDnsykA0vttAUf zrk7*r+zo1DZdLV>5BN`X5ab#J3C!*rd-0+1XTdF)KWL@E(l=xFf{+IlL)i>~QjI}kRr2OAd_sof zV&d04#o(d=>uHN-VcntRIFy`rx~VqaP^7S`W|3RS!g!VK2Q7AMK*A+1i%N4IkDq2n z^v50I<4@xD6d*ipO*-WBp6?5mKlb0^EyROK*TrJ>Mop$oQI@H>Z!~b2ya()y>i1`0tcoX4oj}|99_!!}Q5LSARJ?fGKNeztRhQ571@`lM6oCt_HUS%jxEzZM z(}9irP0bc5ytlYq7={B4bu6L^!3&B4W6W;|f7G9b3B;L}pNRhR)DEHbM|lXreQ7%x z*WGC)SjVSJMvdQ)hW|KA)CaD2#Dj2u?pM3g z2V^B}=je*?#$lr?x%5YQ---+kCx_`ig^I_;)E`YQed!DS$H`l9h92vs4q4sMRBv~7 z*54l`yoZVJ7x)Tjl!mbUP0O#?TY0veVnh)T zi{y0d*5YsVm7?$#OvKMi;P!WFZLQZd32>twPd$W?myUNLqu;ftb0NsUvwy+>X7?yR zBFJ82hQdf*Jauw7TOqRZGBbr2sCMyjfc_=qy&ea~;N1r@{ftBQQ404abl+ph;EAXo ztS>+I&zf8KQ8Zng&3wHI+tvBvK1iPmR75J6p#%awO9TupRC)?T_088cr4bkA3hlTsCQb}^yQvvm;}b(8~ep4 zBw%!{s3+gzQ@o74Y?>6s4?;iUCls(+bmWKh^l;cBD3?}?iY_DjHZK2plArf(68)xO zMDFl1rwLEFw`|p)qS$B)^sfEmOo$6R?r=NQ`i~t?LKgVvPzM2Rv}hfHB;hj(IX@-m z+6M?hh9IMWzvbLQIuw&<{)OJmb@%;cbW42s3rNv_p7sa+aIGNxZo{+Q#-o=R^+$*4 zyY^E`z{~zx+U;Xae!Fy1tuJ^UGx_Me&p2-{G<+FyL4$A8xgVH~6gcBv)gOKRn6=Ou zQ!m8z`0urUvzxYTeY)g0nhrcGY-jt&~W3(h=rC8d+Y ze+&;gmc`)EZDKhs^S6V3Xb^^g-73ah(_{o;;SveKJ;v6|wfFxp+hILXcP6f%BH+|a zSWt#&dUt*4FRLXOW3fSSMiAKXyE;Q5%+~q+C(m(CpU0Kq^Tpg5@dx@!I+`#s@FMQOX!-UWEI>J-Qn0Ep-t&4k5sL?H# zq}iLAyNJ`H-GNdg8M?gSkPwq^?;a1-I-30p!pzOhJ*~E#GH!i)OMf4Olo?W{1@QFz zm*`3IM#`ySPjPkz{L9+lxwiwhvHa1OwY0PxIvC=gDky{j#`xb5=deE8`%jFuc)~`) zJX6;{xw-HR(<`ql5TlD(7CKX5D(IZlt`yaG&oC(@(R&~5l&%dD{*WToGv?My+xzPX zKXaJW`a@-VL5#zZbGJ6mx7Sz@fh@}XCukI8k%In3=ZT2Njyc%_gx_aJj-~^{j^5tA znj3#2_-=-ffQ;NL(oR6%yl~PGKQ9+6Iu63US90>!s6D)c(`OT3tHG^wtu%;;Ne*yU zHsU&efuUMj9-cuge+WDE*iTe?*HaugAXD5{hy}O5Hu5Vw*Bwhb>}0$zb5*;GA#*9LKd8zNZL^4J(gaipgOa0)A1A zvtLEM?h+s_aD46QiA%PvljR5ffeNML+@jj~ySM}dRa49l-F_pA|I+g>;3WS}T2>S# z-R_uI3p^fJe}AAdp5I7JOq_TM`Up`IcFlShShqjpO{rPQFN?oG5eEWFY#i4#i`I(7 zTg(jv(bG2*yD-?zrx5W~5wUBX`|xb(8Z;VS+-$pYJ#q3&c|o7v?}a zW5rY&q2mN}LIcGB$k^|SGV)0U!b$G_;zbC8qJ$YYWp@8EQcww0XDw2}36Hj0pWv|B zMZ~=1`kacxmS69|y*YPr`;UN20OPhjAs+}aQ;JKW(gP@GyhWdia++gm7#t1?0wK4# z{Ds>8@6^UGj!4`FR(0Xbym(1LLJxe4xK4)-efg~dv#Fy#SP*4yJ;K8Y+QvGOMX05S z=hD&9p#d=vMI-t8YmNOZX!c6y<>kGlg;3zPR#i$d1QrO8=%23uaJEiT7>f1lE@U?1VC>CqUDN zpWHER9wX2T3&%TvTsmU{Xz(rW5GiFz0Y6^l)G)sEb-pV6z4hf*;FIHiEoElFi073k zTSZJdZzl@aQHcUb2V^e#{}rB~wRI6t-QaWZnL4T2YV3EK=Q@|D=IQZ1Q!sR&U&KYq z6|T)s#^}P}@pjN)VoYKB_8JSxUrR0Op$g_s=6;26A^~kG2*yCX_UmUD{@lp)U-_60 zx=#^%Bu+!08PKQPj`Q-hOIgX?Mu3ZQ7*^A`B?X1>mbfB0rI=r}gTOGFCk#4iKP`(Q z5<#4OL{6yL=(&%4L+zspL4Rm6#l1kM+V)Eo#)VDb19MgT=&&9(4d75?Ix;#CBkWI< zeRmy~ZMRk%csYnt<5bCDgaj0QfA$*}API@Jn=)TNaJ-KpHjMt{hmNI(WPX)x-~4pxGUp=%!qI%;p;f;%_wqoGkPMMm~>Cpq%1rFY{)NS_)+el|XMlS35ZtGF-?TbLN3iHFJ!%0qK^#r*T z+{k`H{k}j7Ub%TUH~5TmVxANeR;~fD`J?)t3e%n&qHnH4LbG56+Qq-zt>+7{nt(@g z{bilcP|EBjdO=gP8W3lITX4KyuO#HU&L;VWAP)44g{`5jgA5E(Kx(I51>zcCUv|uf zscEx7MT2&3r(>t?VO&ey&H#*T@fvi@DHT<3GZthS+Gpa`7P)W{&%C*v#pTGtfh&6ejWas_ah6MYU71;k zY}ItU$YsgPZDrdLaAMCmSidTXcU-2-B%LzNW#g>|=B@dKQuv)s(Q_hAkN6)(5u&E^ zH%V^aez@Az(GevTLXx=k^LV4VPyFV^j`QOMmyv%e=H9_&9$cI5n4u%Q}(@5XS9^$TsjYF7i$b5s;j-NYAvd?ZR-~)gS7df zNzjY)vr|A-+s)3Oy`ny83GITYPPFIDaoW!67|~@ov~v>WECjG+2wVWJpd+Tgxcx!Z%iK2C=Tt&|K z`v7+$u^tUzPaWo*Dpgg3#*VsK;&x#kM`kyWw^HNczPxZ6%ucs^Fu@Hz&k4;*wgs1KWY3_+}A`2;}kO#d5 z=ZUR>CsLT0GGE`4!O*BaI7O#taW8XRghVHUxelt8vg)b2k3yq3vZ?9VjDd9eIL0dm z1F#PzFK)(YLm0GX?gr%|K)?OFz`uq*1NbN#lNO)`S4prA70JKR%!-_fsCAa-45^cV ze_oIM)>iy=!9X##SxJX;sy%f}d)msHcaXTD5iX{y&Uk9dr@z4fu@U|=?U~#RvUuX; zv#Uj7mN}}{5muG!%OJ0}7U!bbO~&l2Roirz;qqvylV^kyRLFPOuZ+7_FS&a=G<4h< zu5MVk!DRvXXKxLeaEjU&GzA571%PyD&HM-|*RY@2cYTFw@|%Z8?*!y2-+z&IGv>=j zQak1q7q2x=6UpB^+$;H9TOUbjwoU>*E5N6%Mln|=YaO>h_UL4Y*QF)ZbNZ>_42U^E z=k_FGvr)8WQjs1K@+2{COr7;Y^fPDjP;^u}66NMCk0APpe{MW4~jK6lj%ag2|#3EvS6;lAIcLWqJyj4*)0^s0Y1U zz34WWpW0+(B5Q2P{rTk3~kp*XHE?kj`?Wg;h_g^Vzd$tE+Fe z2aJdN&B@5PI>fO?)~aVXz@UxOO-&)Hh0*3SJ~I+v4n#~BmFb`i8dTcH7R*|A5JyFR z+W`q~N9nRrLvwA1bL?@|XO^?)#=FRYFOIv&to0YeEwt;Z6>>|7HVFX&~1O6M7lt4RFy<&4*WBS>F1T)#FXjO7n@sCm?gs@R;w zvC9pu6pvBpI$LAktq#H&*L5#0$|N@aDoN8w*%~)t2Dd% zHGYr-12eo^TmPp9AJ#kljX`i{LAKu5?o%GVWLWbql4WLYT4z}^A)|T-rV?JcvZXJX zSS#?g85_B0u=LLC&&dFF0#RE_s4`9rl28h3k0K~USO{n~qPXXQC2R$V%SnY=+2wwE zyCLWXyPi-db){Pv=3eZ78#J~3CslF9o3yevwC?-jdJZ;YkEFUzGlMBSy0Vr}{rV%f zC`%e1JmOAwtx$gm$Tq-Hd6EOTKR{5rjm(>zP?7uPwU4~}kS=jc|8fVL7~Mo!H2va> zgwPp2r7NxB(e3Ju%j)ung)IZCaYEJ2d*fDH00IbF0}{gKG6-$A7$Iu%ZUHUIr zlbq7UWYY7CwO+i4^Tq!gI9RXWqY(3|`Fy42)6a!-Wuj3>`MRl5vk+e*C`9slIj^Hu!ZhuB%77lWMCn3y3 z%qy}(+Cn^TPt8IaHj~d`ZZKk?w_|u1*y4eX(=gte>qsd9PuKnzF>*^8GuM zzlx3Lm#4~%b#hy&El%q9$F=3bWC9~Ou_L#B(`57H)p-4=4bN`rsrntkn6SGICMg}# zfmv-y*0HAXNkr^2mLxXXs~TrC{{d9mKZRnEEdNFG?2ldWB^9^%P`U6#GFr2JultEQ zX~>4fHom0Q@jA;!X%q$~KPvDzzMixCzlJ(GZ%IBW6Fv~&AmHA}5s?#Og%ZQ)mAJ5VG7Zi@*_c1;1 zPE8|1IY2FGUymDu$KGOGk|jbPWHomz<;=M4Ajj)YWN~%Qny}t*M3vNdL%EQZtAmv@ zHgok=mG#iOp7a{)lcLnHW%6DBP}zj%Wy<6{g!25Ipw}j1VR;eFx0x6uz^;0Wnn|Fp zO~GNiK~0ZNHo9xm!VPqt>pD;Gn{G^@O)FD<9XYgU8%=#E88Hnnr4ePY?s3y1qNK5B z*6LB8-p)?Q`Q2|Fi$~KV?OkD1ck5|dnR1;4zZ`bFX@)vw82{LPV&fDh_WLCqbvbC( z2rwgwEH!Eu5~=!&gzE8X{RagVD{eG7{0dW{qwk|$qk)pjU2D9Ha^#oyFlThZY%=<5 z(UQXpNnhEB8X38aNl)0C-YHIyzTh)jx`nT=DTPgDf()Al7OYnjat@SlWrEG6Si)85 z-3gUh%2cg$7vY;ITKP6jov+d~C3h3S>?;|sdF0`>JUyc%`XPs=UR`vMKNpHVvK<_g3bgaT8 za8<~uqNjH{UrAo|z+6x4E;bnOfim|k?IE_rc46vr$JKN!0MC}U+wl%c0?d2K#X z%$|Fis;YWNsZqU&orK?NxQ7fupg8T}EsgqJ!^BQgLpIb)iLlz6eTHI;RSLy+beE_{ zI@iSCG-K$~qrz(qCK*G7kIR+EV~fV_SIR$I!7>MJJQNr{k&d8ZzkKD|3%Bs3ccx}$ zo|6l~-QDs;So~GUaTlAg7)c)S_?=P8LSpW8pC{p7BJjO{H})yt@0v_hu*;4-k$$?E zL`GQY5V@+mYu2H?05bK#_=L%5QnsPnGo~;oy42&-s9dxZyfB@@z6sD8$*9_r1zg0n zZrcy-dJNNgk^qGP3 zGZT1K2gliXIu_-b_yln9f@KSf-`bFCqI91 zOY=;|#vgU6I;!g!hJ}MOv3^^X702!vaMpwuEKF8JCx`(1gYdEkEjxO*h<;1^$cBNz zfTzY@Wo&q9%(qyl<0x6{lH5Iu0#v1jP!23x9I|0+4DBuou_-&|Mq6W4f8iK;e#?_0 zckX*Gv7^?p1e_z*v0?`0%10il_U-O3TR;$ojg^^%O!^$JIUJ(nlZCGjr+tj~_G_^u zVf(*eZ4hpaM3?6q?tBIhrL=%2ZJohtDh#v-#5`2e8zYWdze4J?2QsLmi8g)X3tkxn zGl-9crdue*s>UUPOc!gqew};leh;sEr}UDTcet2_K}PZL{V6_Y6PJ~_{;i1IW9Z#sLRv-4SSijV-lNwBN&>9hz4Mqmg$qv&#KmVa_11Q@}_)a z25)*4S$<~Ls$0E~X_C}{b+l<5Z+)f~qutp3sbVcw?Z}mX$`0#5pi!>P3XHjj!HT zQz8qWyZr*Omwx@*Ak)Z-Vp$%!@rBG&#UzV?GbW|6`)la~*0G>JZ=mK~uAIVDzyzAr z0p>IptrwpCK8u#E3S3jgdt4U2nq8i^JoAn9hsYAkMReym67H|W@5D-4f$&Z&b^f&1 z)?L>j=FM9M$Ecy(Yr^Ld1$jN~Pc}*hh1Yn=QLfV<*?6Xd_^=Z$)>y7ouI!PWKE2rv z@o$&Jo1UssJ&pR)1M|i$LChfr9fx%R>6VEuMebt?Xmk##;TVs9x`%smbU8@pEYsNI zH&Tf7m|U>>RFD})$;3Bi)t6sZz)wl!Oa<&mD9X3ix3)t%SF$!6{4WIrOv=7Bu;(cC z0_X@_8viQQc}-LQ4r2WXuz^ke6O)wRZvzHNcIBh?G`?(3EVI?N$?B)VtwW7L2IHUP zWv+3acDZGX`SiXpGv;yOXL(A!kx4I8P47GCgav`v7dR)76&?bxt^TF(({Rdelv z#IDgKNQHmM`!IQvEk$Mu@lL>@B>aR&-I-y8`K{UiEaBYnz76O2jOo&mqh$d5_VO_y z^{;7Ar5(YU!|P$4`weWvOpLM$w(=ovFg|DbtxkLeemS#QpguG7Dc`C>@mNliGWFMl z05V4MN-nyO#u;pUB)vpnil?GIB$k0rFZi;lP%H0Wz!va2!HcNqqz@8wX{t{^&wqc_wiOx7@jE>*qvFGt*!>X( zgZmq&bUE_+$!B4q38kv@z2{#+Q93ps4v#wtg(j-`%rvS!>~iZ`pR_5#_a4WD!IdE} zBAfOkHkgfAhow%{EBC4j*NA?Z@>Mpdb*cDP=rn%Hu+vL>-SKXTh)*o@4T%+lhZPOg zONcn;4ey!)dHqDa`^8G(@ulSpPk+Cx7WIIsebJ?fo7jk3zI>V~bn1v>$Lt#(gh$QZ zo_lj(D+Rfvde364R|mpkc$~xtmR1P$Kpl{WG+q+q1C?B&*)ke$tmF7gtmx{+>+-AL z1WI;!;>vKw?tckHc9CPCwz^wYwx%m8w+`c{8&NF$XHVH=m7eYjY`_$du3)!EZ4t=1 zEZc*Dx3uiTxlaeA9e_k_o{YYB0+8ZWooPnJbIO~Vf9%Hj6EsV}t&RtXVlmNB^vmY@ z{AyHShKsx1?_b=9D|3YIOXar@7rNWT++k`q&5~@M>G8|6V2+!8IO&^ptTqWy9s9MR z&KB@Q$=^VRQ&<8Sj?uYN5e#lYU2WdQ06>9*S&2CWz(vePVY2rtR;BV>+>84P;7daM zGDk{dv_uPv+so`~WS()Fgc~u-l3^R{!;?7*5YDK2w1y8N%6t;=Hp|}Jv7EE#bfFg6 z;G5PM>)U?@OnoF4HSba?&0h3}d3s+HBDULG)!d6$Y77ro4u~Lj>f-o^K@j7Pv@&n9u1I-Ozq-RDEYTM2bW8RKns+J zf(#?4ccldUdwk?=hG3`zi{w#8>dh>y9HYy53T;30|if$sY@m)+*md&X5DS*n z_0ff;ZN3mx@n``_h~Q1o>3@_jB{M*gbXYsrI5FeiQdXCrooyev43v;&JMLz?!R}1I zs+jxMWet!;9lbisI|J$)>?kt&9(oID+2fRo8Cf2M@pI_Vrz$`$J}lmarWHvnKm~u! z&G_EX_2unw%B963pRO_Ylv;i3Ju()TzQRCyD}I zKn0z?{shXfr==HzH{g{BqJLqtFM4{%MKw3s!f7kdVhm%)8tn!9oEYAA^v*Mc3dr(VotJVrd8xM=C}xSj}J*F3Wh zt(A-n8NzOe+ow9DpP6T`H2!#RjA&d9??VzuD$Ir-#_p%S$p)8*Y{xBdO2=&xV2|r* zU?+mb<$(2#vHK*B7t<5z%KxGfr=UO=5?j5c4{T%+UHmm@w>f5O5eXdisXQphQp2EC z^d%H4Cai1-(f-j;5sX$~uibQ3;=HbJJ~>1=Ic(tlQDnW# z4_)OFwWU5?zK_JyeJB_G>5R5kxw-;)j~$ot&kPX%C6dC9%-fylBE5Q~!RrjMYBCuL z{4B){kwiglDGA`c&X$^dyKpYYbI>zaN=D3MRb8Akt+9Hjrq+?(q3x02-Ob3@FS>)v zji>@H<;|E_Z@?<3-UdljDjFb?eBvE8qk=d4!&8_uOLT?P0`Qiw#qmf`+Sxo~0nN-_a&2WgmZvloC6eVc4etpA_Eh$$T-oyZJG77u z?LNRTj8G<00hx%zPV0Hnus`H1nE3DfRT(#8KY`flb9=7cvP_Y0=s zv(4Hx&#k_6MI}(dOLFZSQFniR>(Vd^N#ddf1FnPn%L=+o#{ng^_Lw@ndOv+5Zk-AqikLIlix zH6m$JCr@ngKQ&x#Y1V&bi6IcfKYjd5E;6dTxo+@VfYJBSH%`RpT9DGP5~J~A-$?wj zjhk&e`SOld{hp)j)Irj1>slEWfr`-TSYhBI8M;oHP!9T00xp+Ifw>So0Ob%U%&cR> zW9{R@41fNapt`B~LxIBED(YS1Gs_>juDiKG*`r&j(7-2XaDpku1LZ7{@8gPZ>tgqv z$R8fR^J;WRTzz{OpAzn=-O2B2tu_r$Y%?>~QI7_(O7e~KK?$@MRC*EcfJ@Vvyet7e z4nEFJ0PY?ez30cR+?X#nN$(Y3&pXtmR^d|jY$=04iKo?_k^-iHVzlTWE1#f-*mb$c zuU|RQ$?U%IFKzTPxvQMrU~y1!r8*b%5Texol`lXhv;r#HXuFtI1kRNN*G16_CV141 zgAqX+A^@8s?-`+<=>+7WUvjQ5D+K0mO=lS4>+Q>hu_IlZg_!Y23gY#27`v7Zrr}h+ z)k`6k7Iw@%v|8oLNqTp03U@**^2+wk;gJwTB5OW`#^S&e!4RcSZ8x$-gISZmdYswg z&}3!3)0Q4u)o_H5ZcT>_!i)#XLsv~2A(ug)UWCa(GbFZWYo4B6Hi%{d-MWcB%w6tHJ-Ssk(JO=2b z&R3mQlEfJH>`BAwG`~ujYa_dR>A96irWXjvb6*@AKP#5Yj?ZR}F|4@{Of%VmTJ8OS zJ`0dns*guI^~Ky4NaR15z@Tt$Obg4*?J4^73=Y=~3-{ug26$q_QGj*04!cFYSX}f} zUvhU$-!vQ*W@et}Itt>@MwQBq`@kibVu8{<}|n?lN|1B^zt- z&Pb6i6p?Pk3i)4E0Lfr9+m>D35QB894BW9yBi>p5P<2u z05G=uLx)W(kupg_!k4dSr7H{9Curb0Pz=yq$Vx^ zmnq->Lf`#hEGLW%yj#=u!tg))k8>q?4_t9R7RHEZ+(mZ#cqbX3AKEgX4M~l$ZxJ%z z-8Sm*S+z|(?$yX@Agw!=aJTy=v4L)U4(p;vNcz*$Z$$IFX}aC@-c?~lwpGYe$jvI`96*MJnVL$*Npe8gF|-5{=S|RQ6+R^A03Oao?%udWjJymBRx^< z{@RQ3J5B`w2=3n0e)O zr{?6|^5o-df-Y6sBwy5$qSKK3rrvz>FWrKV>=G;H32I$C2E*5X!p;q;I3t=`MK+#% z=e*mdBL}jY);mCbJ^*=$KgzRH+JXqGub83TGK+eabu<4Hf9P^qLOSHoU;-eru+uZY zxWiCuW3OB_1kh&*XZP7R#`YC~rJ1ZE{Cl=jHnI;C?3p^BwIMM1a_oI>dbZQGwQe=HJSTri>x__|z+({iq?4Y~dS z11Gt-r8p16ml-)8OWQnUY$EcC^j`PkaQ9`^x_!HjS$!{FHf{M9&ZweBPB{y(9OVkQ z@Qv&s!{DzJ`Id)3k>r}^pIvQzo0R#_TM6J5A?1GIGrTX?T&=Y?PAh69OR{pFardgh zzE$C)#CMa*rhW>n$9grrNL&qaR|>A0UT~6`@qc^1DXXj16?jWV9guM0ZlG5|aBmMF z1B@>K;QB+CDDGWITfhFiN`Z-gH`GG=#`P#L^*(-GXPq95MZvTr8tAF3zZNXwE%9Xx zJ$~lR>zQk<%k#{^Xv0BNPCEIU${Se)#Q?U4H&@GB9IBUo8v=KJU<58NW_cI{|9%kx zc+ap$NoKZee+6#1!is2}+*!d=nMD`#p|NxAW*YqQ?&?7Xx8=xanXQF*NAOXY)vO{V zlcROTEmh0$T*%(W?@u#1n5VX?rsuo|{WW*mUgHe8yw1O06anBS-~GUJMLSX0GxA#6 z87bL=71_g`km+bb;p=f6V8yN1;VJ^nZFe_i(7zJ&x~ODyNWo zPD*VlwN*$pq%he^YGNpr+{P%EO}T_E%@`)BOcx1-+@?fQ6Eessm&zrBgl!y!$^9~B zCdO@;%lR$sv!Ap7={bL$^@rz~#WU-Dz3=;*HS7I;K2Fvz#1A5qj8uakWSVzxL>guKVDV_c=e%n%L3HhZlwf|!NbOg% z!eDMVhjNPwNWwr6Tg3WxmDW*IG}QnHbZV3KEL**-Nt(A=%-Q8YfB5`rE(5#zTK6Vu zZn;Y3<*=OvNVT)Itw@EN`k`F#HQ>%hPBAE#Lu#{$Z#LMgH z@1uXzO$)lqLI;2p^vSOz8h=w+TOwLS@a%B2=8a6JBnLh#c}dOARve8EOH4O<4tJ}# zLFTE!TmRMK(dz%;VALX$OU;K=F#!Pqk#ekf%k&HDW$Din-<0*px&=SHo_4M?XK{LC zw9-LVT!v@WyBPmH?{2Md44{ZVR0og(lI#fj=@;u;R3 zdHVsgSGcO?_FA{Go#lcw{*lBTxXlIzZ3BW%#ois;x9<-3$SsDwNTR4KUd!g9uoQl- zOcMBj1)y9i=D!BSB3Zd8TG+)8*Fr)8<0F`(M>qYVYvtoZ6HO#m+|$l`V+%poecQHf z)c|xxZl{OxL?qKLee2B z$gMD^@2`ZtH#OU=lQ*}faPp~+Q*pyUIF{>#kop1RdfM9OX97OSe>JQSMsL>AvUGGz z{@!*nIMF8X(kVQs%Wer@^UFB}_pGgw6Z6XFnCtOqYlNvxH447D&=J+SigCQCSGC~M z_}RpHv|_K{ZK-@Ya1Il&%f6B#SPrZ#OS2c_suDlw*8d@PaPbOl;%#Z4spNf5PI8j| zh7?2pJsymzD)ad0mXu%~V943ii;+hNUb`m!9ys|i4uRsWbAN5C2Ct^%NEb=a_@qa@ zO2qGIlO~_vfIr(pB8jKkm9m0xpKuB~@Y1VeIxEqoKSBBe0&@NytSTUzsSXkaZ(9ingC^Ueo_Q7hz}5?X7%pM0S5?PJbK#J&aTGv z7}D5TdSz+f9b*tsRbUlp$XdS0{p=xk%T$2|5|jl53c0AWt30@@DkEfZAV+(5kZfo6 z1S0-S@K0>btB{mBU+nNeIM(`fbZ9>;zBujNk^|S5uRXiU-jGeA6S_~G#%)nXPu@-4 z8}M#A{rK_Y6CAwMv3$OZ&!ji)*9P?ky^IG{B-T33s1A%LdVtU>QOV%3nyAN1S|~YM zRFW;^;~Vqv(pbDB7jmMA+btgcOr-~=Nv`RBN1-^hRC|E6z z4lco9FuOAPS}gSpG7evhlcXb>LU(2%sXfb5RiF+K@@^T-wY$8~dG4zLu!8_#9)`2G zPf7)6Ne_L)zCd%NfYw7b+yZs0`%<4G;cLVapLsoirR$i?z)R8DY!4|Kjc#h^k8#nH zq1co?jbrP8yiuVPq3l7xn1IYk+hfNxz-)JagLwEbB3Y62EHr+*I|s1ZQizNh+jDy{oDxM_RKWrspqmq~Z$cZ&#y@BL+r7I>tTVcs>{@+I=V$Um(+k^z`FW&mA083^Qxz@Ee=38F^=Zc->avn;r8TP%XK!B+*7Kri0} znFP*>Gf%|=vZ$7vl=72)Xj@rm2uK#~y;oo8enSwwmq&VAp`gQ!w^i}*Z+$4XmI_Rz zZgN*{X3j0uX)HGpPqtdjnn-5*l&~BQ>qR`!N?o1zvzrk07NHZYXLdCEjjy_i&k0?R zWA>k`cVVouPlq*R`}FmYa4wBAp Date: Tue, 22 Sep 2020 23:40:43 -0400 Subject: [PATCH 08/10] update images --- README.md | 6 +++--- img/{Scan runtime.png => scan-runtime.png} | Bin ...on runtime.png => stream-compaction-runtime.png} | Bin 3 files changed, 3 insertions(+), 3 deletions(-) rename img/{Scan runtime.png => scan-runtime.png} (100%) rename img/{Stream compaction runtime.png => stream-compaction-runtime.png} (100%) diff --git a/README.md b/README.md index 6922f9c..029c671 100644 --- a/README.md +++ b/README.md @@ -8,7 +8,7 @@ CUDA Stream Compaction * Tested on: Windows 10, i7-8750H @ 2.20GHz 22GB, GTX 1070 # Scan, Stream Compaction and Radix Sort -![].(img/main.PNG) +![].(img/main.png) ## Performance Analysis @@ -16,11 +16,11 @@ The runtime of the naive method on GPU is worst, followed by that of CPU, then G Both the naive and work-efficient method on the GPU only rely on global memory instead of ultilizing shared memory, which potentially lead to much slower runtime than thrust method (and CPU for the naive method.) Naive and work-efficient methods also do not use up the potential of warp partitioning or memory coalescing, which are features that thrust may have to reduce its runtime. -![](img/Scan runtime.PNG) +![](img/scan-runtime.png) The same explanation above is applicable to the below as well. -![](img/Stream compaction runtime.PNG) +![](img/stream-compaction-runtime.png) Commparisons of GPU Scan implementations: \ No newline at end of file diff --git a/img/Scan runtime.png b/img/scan-runtime.png similarity index 100% rename from img/Scan runtime.png rename to img/scan-runtime.png diff --git a/img/Stream compaction runtime.png b/img/stream-compaction-runtime.png similarity index 100% rename from img/Stream compaction runtime.png rename to img/stream-compaction-runtime.png From 3411b987f6e7152572409be4e70e66c36eacc025 Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Tue, 22 Sep 2020 23:41:49 -0400 Subject: [PATCH 09/10] updated cover --- README.md | 2 +- img/{main.png => cover.png} | Bin 2 files changed, 1 insertion(+), 1 deletion(-) rename img/{main.png => cover.png} (100%) diff --git a/README.md b/README.md index 029c671..00c6d05 100644 --- a/README.md +++ b/README.md @@ -8,7 +8,7 @@ CUDA Stream Compaction * Tested on: Windows 10, i7-8750H @ 2.20GHz 22GB, GTX 1070 # Scan, Stream Compaction and Radix Sort -![].(img/main.png) +![](img/cover.png) ## Performance Analysis diff --git a/img/main.png b/img/cover.png similarity index 100% rename from img/main.png rename to img/cover.png From a9da2d43473d999f722a0fbe3a688a314bb45cb5 Mon Sep 17 00:00:00 2001 From: "Thy Tran (Tea)" Date: Tue, 22 Sep 2020 23:57:00 -0400 Subject: [PATCH 10/10] updated readme --- README.md | 77 ++++++++++++++++++++++++++++++++++++++++++++++++++++++- 1 file changed, 76 insertions(+), 1 deletion(-) diff --git a/README.md b/README.md index 00c6d05..c196685 100644 --- a/README.md +++ b/README.md @@ -10,6 +10,8 @@ CUDA Stream Compaction # Scan, Stream Compaction and Radix Sort ![](img/cover.png) +The main goal of the project is to implement different methods for scan (mostly exclusive scan) and stream compaction. Scan is used to aid with stream compaction and radix sort. For scan and stream compaction, several CPU and GPU methods are implemented to compare their efficiency. + ## Performance Analysis The runtime of the naive method on GPU is worst, followed by that of CPU, then GPU work-efficient method, and finally by thrust. @@ -22,5 +24,78 @@ The same explanation above is applicable to the below as well. ![](img/stream-compaction-runtime.png) +Another thing to note that greatly affects performance within GPU implementation for both naive and work-efficient methods, compared with CPU method, is that the GPU ones only work well for a power-of-2-elements array. If an input array is not a power of 2, the number of elements is increased to the nearest power of 2. Hence, in this case, the runtime for scan and stream compaction on GPU is not that much better than that on the CPU. + + +The below is an example of the results with 1000000 elements. Note that the CPU scan time is 0ms (which is not true) because the CPU timer is disabled here, so that the CPU timer for stream compaction that uses CPU scan does not stop the program from executing. + +``` + +**************** +** SCAN TESTS ** +**************** + [ 2 30 36 0 31 22 46 36 16 3 45 48 20 ... 37 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 0ms (std::chrono Measured) + [ 0 2 32 68 68 99 121 167 203 219 222 267 315 ... 24482173 24482210 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 0ms (std::chrono Measured) + [ 0 2 32 68 68 99 121 167 203 219 222 267 315 ... 24482069 24482112 ] + passed +==== naive scan, power-of-two ==== +kernel invoke count: 20 + elapsed time: 2.18784ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== +kernel invoke count: 20 + elapsed time: 2.35523ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 1.20525ms (CUDA Measured) + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 1.19664ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 1.94966ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 1.9712ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 1 0 0 0 2 3 2 1 0 3 1 2 1 ... 0 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 2.3196ms (std::chrono Measured) + [ 1 2 3 2 1 3 1 2 1 3 1 2 2 ... 2 2 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 2.3835ms (std::chrono Measured) + [ 1 2 3 2 1 3 1 2 1 3 1 2 2 ... 2 2 ] + passed +==== cpu compact with scan ==== + elapsed time: 8.1008ms (std::chrono Measured) + [ 1 2 3 2 1 3 1 2 1 3 1 2 2 ... 2 2 ] + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 1.1991ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 1.19808ms (CUDA Measured) + passed + +***************************** +** RADIX SORT TESTS ** +***************************** + [ 5 10 32 46 46 31 8 9 28 23 47 42 41 ... 8 0 ] +largest number of bits: 6 +Correctly sorted: + [ 0 0 0 0 0 0 0 0 0 0 0 0 0 ... 49 49 ] +Yours: + [ 0 0 0 0 0 0 0 0 0 0 0 0 0 ... 49 49 ] + passed +Press any key to continue . . . -Commparisons of GPU Scan implementations: \ No newline at end of file +``` \ No newline at end of file