From 5e71147348e57c6ddc4da12406c7e644acaa6be3 Mon Sep 17 00:00:00 2001 From: Nadine Adnane Date: Tue, 17 Sep 2024 23:44:43 -0400 Subject: [PATCH 1/3] WIP --- README.md | 29 ++++++++++--- stream_compaction/cpu.cu | 69 ++++++++++++++++++++++++++++--- stream_compaction/naive.cu | 83 +++++++++++++++++++++++++++++++++++++- 3 files changed, 168 insertions(+), 13 deletions(-) diff --git a/README.md b/README.md index 0e38ddb..2e10cd6 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,29 @@ 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) +* Nadine Adnane + * [LinkedIn](https://www.linkedin.com/in/nadnane/) +* Tested on my personal laptop (ASUS ROG Zephyrus M16): +* **OS:** Windows 11 +* **Processor:** 12th Gen Intel(R) Core(TM) i9-12900H, 2500 Mhz, 14 Core(s), 20 Logical Processor(s) +* **GPU:** NVIDIA GeForce RTX 3070 Ti Laptop GPU -### (TODO: Your README) +### Results -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +Note: I am using a late day on this assignment. +## Part 1: CPU Scan & Stream Compaction + +## Part 2: Naive GPU Scan Algorithm + +## Part 3: Work-Efficient GPU Scan & Stream Compaction + # 3.1. Scan + # 3.2. Stream Compaction + +## Part 4: Using Thrust's Implementation + +## Part 7: Write-up +# Comparing Implementations +# Performance Analysis +# Test Program Output +``` ``` \ No newline at end of file diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..932d201 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -19,7 +19,14 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + // DONE + // Exclusive - Include the identity in the output + odata[0] = 0; + for(int k = 1; k < n; k++) + { + odata[k] = odata[k - 1] + idata[k]; + } + timer().endCpuTimer(); } @@ -30,9 +37,21 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + + // DONE + int j = 0; + for (int i = 0; i < n; ++i) + { + if (idata[i] != 0) + { + odata[j] = idata[i]; + ++j; + } + } + + timer().endCpuTimer(); - return -1; + return j; } /** @@ -42,9 +61,49 @@ namespace StreamCompaction { */ int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + // DONE + + int numRemaining; + + // Step 1: Compute temporary array containing + // either 1 or 0, depending on if element meets criteria + int* temp = new int[n]; + for (int i = 0; i < n; ++i) + { + if (idata[i] == 0) + { + temp[i] = 0; + } + else + { + temp[i] = 1; + } + } + + + // Step 2: Run exclusive scan on the temp array + // Exclusive - Insert the identity + odata[0] = 0; + // Start at 1 since we inserted the identity and are shifting to the right + for (int k = 1; k < n; ++k) + { + odata[k] = odata[k - 1] + temp[k - 1]; + } + + // Step 3: Scatter! + // Result of scan is index into the final array + numRemaining = odata[n - 1]; + for (int i = 0; i < n; ++i) + { + if (temp[i] == 1) + { + odata[odata[i]] = idata[i]; + } + } + timer().endCpuTimer(); - return -1; + + return numRemaining; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..015d0c7 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -11,15 +11,94 @@ namespace StreamCompaction { static PerformanceTimer timer; return timer; } - // TODO: __global__ + + int* dev_idata; // input data + int* dev_odata; // output data + int* dev_tdata; // temp data + + // Define the blockSize + int blockSize = 128; + + // DONE: __global__ + __global__ void naiveParallelScanKernel(int n, int *odata, const int *idata, const int offset) + { + // Using the Naive algorithm from GPU Gems 3, Section 39.2.1. + int idx = (blockIdx.x * blockDim.x) + threadIdx.x; + + if(idx >= n) + { + return; + } + + if (idx >= offset) + { + odata[idx] = idata[idx - offset] + idata[idx]; + } + else + { + odata[idx] = idata[idx]; + } + } + + // Insert the identity and shift right + __global__ void makeExclusive(int n, int* odata, int* idata) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + if (index == 0) { + odata[index] = 0; + } + else { + odata[index] = idata[index - 1]; + } + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + // DONE + + // Memory allocation for input data + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + checkCUDAErrorFn("cudaMalloc dev_idata failed!"); + + // Memory allocation for output data + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + checkCUDAErrorFn("cudaMalloc dev_odata failed!"); + + // Copy from host input data to device input data + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAErrorFn("cudaMemcpy to device failed!"); + + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + + // ilog2ceil(x): computes the ceiling of log2(x) as an integer + int numLevels = ilog2ceil(n); + int offset; + timer().startGpuTimer(); - // TODO + + for (int d = 1; d <= numLevels; ++d) + { + offset = 1 << (d - 1); + naiveParallelScanKernel<<>> (n, offset, dev_odata, dev_idata); + std::swap(dev_idata, dev_odata); + } + + // Insert the identity and shift to the right to make exclusive + makeExclusive<<>> (n, dev_odata, dev_idata); + timer().endGpuTimer(); + + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + checkCUDAErrorFn("cudaMempcy to host failed!"); + + // Free the dev data + cudaFree(dev_idata); + cudaFree(dev_odata); + checkCUDAErrorFn("cudaFree failed!"); } } } From 28d316d74bacb77f51e71483d782a2e95740d706 Mon Sep 17 00:00:00 2001 From: Nadine Adnane Date: Wed, 18 Sep 2024 22:43:34 -0400 Subject: [PATCH 2/3] CPU implementation bug fix, debugged efficient implementation Also added test output to README --- README.md | 57 ++++++++- src/main.cpp | 2 +- stream_compaction/common.cu | 24 +++- stream_compaction/common.h | 1 + stream_compaction/cpu.cu | 2 +- stream_compaction/efficient.cu | 203 ++++++++++++++++++++++++++++++++- stream_compaction/naive.cu | 102 +++++++---------- stream_compaction/thrust.cu | 19 ++- 8 files changed, 337 insertions(+), 73 deletions(-) diff --git a/README.md b/README.md index 2e10cd6..381fb9a 100644 --- a/README.md +++ b/README.md @@ -27,5 +27,60 @@ Note: I am using a late day on this assignment. ## Part 7: Write-up # Comparing Implementations # Performance Analysis + + + + # Test Program Output -``` ``` \ No newline at end of file +```**************** +** SCAN TESTS ** +**************** + [ 3 7 33 47 27 3 16 21 20 14 40 30 20 ... 29 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 13.8246ms (std::chrono Measured) + [ 0 3 10 43 90 117 120 136 157 177 191 231 261 ... 102740616 102740645 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 13.7131ms (std::chrono Measured) + [ 0 3 10 43 90 117 120 136 157 177 191 231 261 ... 102740526 102740542 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 4.41216ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== + elapsed time: 4.15386ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 2.04285ms (CUDA Measured) + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 1.68448ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 76.0596ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 46.633ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 2 3 1 2 1 0 1 1 2 2 0 2 3 ... 3 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 26.062ms (std::chrono Measured) + [ 2 3 1 2 1 1 1 2 2 2 3 1 1 ... 3 3 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 26.0387ms (std::chrono Measured) + [ 2 3 1 2 1 1 1 2 2 2 3 1 1 ... 2 3 ] + passed +==== cpu compact with scan ==== + elapsed time: 71.3503ms (std::chrono Measured) + [ 2 3 1 2 1 1 1 2 2 2 3 1 1 ... 3 3 ] + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 2.31478ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 2.08794ms (CUDA Measured) + passed``` \ No newline at end of file diff --git a/src/main.cpp b/src/main.cpp index 896ac2b..ce9bb93 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1 << 22; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..379cdb9 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -23,7 +23,16 @@ namespace StreamCompaction { * which map to 0 will be removed, and elements which map to 1 will be kept. */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { - // TODO + // DONE + int idx = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (idx >= n) + { + return; + } + + // Map non-zero values to 1, otherwise 0 + bools[idx] = idata[idx] != 0 ? 1 : 0; } /** @@ -32,7 +41,18 @@ namespace StreamCompaction { */ __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { - // TODO + // DONE + int idx = (blockIdx.x * blockDim.x) + threadIdx.x; + + if (idx >= n) + { + return; + } + + if (bools[idx]) + { + odata[indices[idx]] = idata[idx]; + } } } diff --git a/stream_compaction/common.h b/stream_compaction/common.h index d2c1fed..b2b34dd 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -17,6 +17,7 @@ * Check for CUDA errors; print and exit if there was a problem. */ void checkCUDAErrorFn(const char *msg, const char *file = NULL, int line = -1); +const int blockSize = 128; inline int ilog2(int x) { int lg = 0; diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 932d201..bffa308 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -24,7 +24,7 @@ namespace StreamCompaction { odata[0] = 0; for(int k = 1; k < n; k++) { - odata[k] = odata[k - 1] + idata[k]; + odata[k] = odata[k - 1] + idata[k - 1]; } timer().endCpuTimer(); diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..c7f75fe 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,13 +12,115 @@ namespace StreamCompaction { return timer; } + // Helper functions + // Up-Sweep - parallel reduction + __global__ void upSweep(int n, int d, int* data) + { + int idx = threadIdx.x + (blockIdx.x * blockDim.x); + int index = idx << (d + 1); + + if (index + (1 << (d + 1)) - 1 < n) + { + // From the slides: + // x[k + 2d+1 – 1] += x[k + 2d – 1]; + data[index + (1 << (d + 1)) - 1] += data[index + (1 << d) - 1]; + } + } + + // Down-Sweep - traverse back down the tree using partial sums to build the scan in place. + __global__ void downSweep(int n, int d, int* data) + { + int idx = threadIdx.x + blockIdx.x * blockDim.x; + int index = idx << (d + 1); + + if (index + (1 << (d + 1)) - 1 < n) + { + // Save left child + int temp = data[index + (1 << d) - 1]; + + // Set left child to this node’s value + data[index + (1 << d) - 1] = data[index + (1 << (d + 1)) - 1]; + + // Set right child to old left value + this node’s value + data[index + (1 << (d + 1)) - 1] += temp; + } + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + // DONE + // This approach uses balanced trees to avoid the extra factor of log2n work + // performed by the naive algorithm. + // GPU Gems 3, Chapter 39.2.2 + + int* dev_idata; + // Calculate number of levels needed for the scan + // ilog2ceil(x): computes the ceiling of log2(x), as an integer. + int numLevels = ilog2ceil(n); + // Calculate the power of 2 number of levels + int numLevelsPow2 = 1 << numLevels; + + // Memory allocation + cudaMalloc((void**)&dev_idata, numLevelsPow2 * sizeof(int)); + checkCUDAError("cudaMalloc dev_idata failed!"); + + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy dev_idata failed!"); + + int padding = numLevelsPow2 - n; + if (padding >= 0) + { + cudaMemset(&dev_idata[n], 0, padding * sizeof(int)); + checkCUDAError("cudaMemset dev_idata failed!"); + } + timer().startGpuTimer(); - // TODO + + //================================================================================ + // Part 1 - Upsweep phase + //================================================================================ + for (int offset = 0; offset < numLevels - 1; offset++) + { + // Calculate the number of blocks + int numBlocks = (numLevelsPow2 / (1 << (offset + 1)) + blockSize - 1) / blockSize; + + // Perform the upsweep phase + upSweep << > > (numLevelsPow2, offset, dev_idata); + checkCUDAError("upSweep kernel failed!"); + + // Sync before proceeding to the next iteration + cudaDeviceSynchronize(); + } + + // Need to set the last element to 0 before starting the down sweep phase + cudaMemset(dev_idata + numLevelsPow2 - 1, 0, sizeof(int)); + checkCUDAError("cudaMemset dev_idata failed!"); + + //================================================================================ + // Part 2 - Downsweep phase + //================================================================================ + for (int offset = numLevels - 1; offset >= 0; offset--) { + // Calculate the number of blocks + int numBlocks = (numLevelsPow2 / (1 << (offset + 1)) + blockSize - 1) / blockSize; + + // Perform the downsweep phase + downSweep << > > (numLevelsPow2, offset, dev_idata); + checkCUDAError("downSweep kernel failed!"); + + // Sync before proceeding to the next iteration + cudaDeviceSynchronize(); + } + timer().endGpuTimer(); + + // Copy the results back to the host + cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy device to host (dev_idata to odata) failed!"); + + // Free device memory + cudaFree(dev_idata); } /** @@ -31,10 +133,105 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { + // This stream compaction method will remove 0s from an array of ints. + // Initialize necessary buffers + int* dev_idata; + int* dev_odata; + int* dev_booleans; + int* dev_indices; + + // ilog2ceil(x): computes the ceiling of log2(x), as an integer. + int numLevels = ilog2ceil(n); + size_t paddedSize = (size_t) 1 << numLevels; + + // Allocate memory for device arrays, copy input data to device + cudaMalloc((void**)&dev_idata, paddedSize * sizeof(int)); + checkCUDAError("cudaMalloc dev_idata failed"); + cudaMalloc((void**)&dev_odata, paddedSize * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed"); + cudaMalloc((void**)&dev_booleans, paddedSize * sizeof(int)); + checkCUDAError("cudaMalloc dev_booleans failed"); + cudaMalloc((void**)&dev_indices, paddedSize * sizeof(int)); + checkCUDAError("cudaMalloc dev_indices failed"); + cudaMemcpy(dev_idata, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy dev_idata failed"); + + if (paddedSize > n) { + cudaMemset(dev_idata + n, 0, (paddedSize - n) * sizeof(int)); + checkCUDAError("cudaMemset dev_idata failed!"); + } + + dim3 fullBlocksPerGrid((paddedSize + blockSize - 1) / blockSize); + timer().startGpuTimer(); - // TODO + + // Map to boolean + StreamCompaction::Common::kernMapToBoolean << > > (paddedSize, dev_booleans, dev_idata); + checkCUDAError("kernMapToBoolean failed!"); + + // Perform scan on the boolean array + cudaMemcpy(dev_indices, dev_booleans, sizeof(int) * paddedSize, cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy from dev_booleans to dev_indices (Device To Device) failed!"); + + //================================================================================ + // Part 1 - Upsweep phase + //================================================================================ + for (int offset = 0; offset < numLevels - 1; offset++) { + + // Calculate necessary number of blocks + int numBlocks = (paddedSize / (1 << (offset + 1)) + blockSize - 1) / blockSize; + + if (numBlocks > 0) + { + upSweep << > > (paddedSize, offset, dev_indices); + checkCUDAError("upSweep kernel failed!"); + cudaDeviceSynchronize(); + } + } + + // Need to set the last element to 0 before starting the down sweep phase + cudaMemset(dev_indices + paddedSize - 1, 0, sizeof(int)); + checkCUDAError("cudaMemset failed!"); + + //================================================================================ + // Part 2 - Downsweep phase + //================================================================================ + for (int offset = numLevels - 1; offset >= 0; offset--) { + int numBlocks = (paddedSize / (1 << (offset + 1)) + blockSize - 1) / blockSize; + if (numBlocks > 0) + { + downSweep << > > (paddedSize, offset, dev_indices); + checkCUDAError("downSweep kernel failed!"); + cudaDeviceSynchronize(); + } + } + + //================================================================================ + // Part 3: Scatter + //================================================================================ + StreamCompaction::Common::kernScatter << > > (paddedSize, dev_odata, dev_idata, dev_booleans, dev_indices); + checkCUDAError("kernScatter failed!"); + timer().endGpuTimer(); - return -1; + + //================================================================================ + // Part 4: Copy results and free memory + //================================================================================ + int numRemaining; + cudaMemcpy(&numRemaining, dev_indices + paddedSize - 1, sizeof(int), cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy numRemaining failed!"); + + // Copy result from device to host + cudaMemcpy(odata, dev_odata, sizeof(int) * numRemaining, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy odata failed!"); + + // Free memory of temp buffers + cudaFree(dev_booleans); + cudaFree(dev_indices); + cudaFree(dev_idata); + cudaFree(dev_odata); + + return numRemaining; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 015d0c7..9676797 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -1,5 +1,6 @@ #include #include +#include #include "common.h" #include "naive.h" @@ -12,93 +13,68 @@ namespace StreamCompaction { return timer; } - int* dev_idata; // input data - int* dev_odata; // output data - int* dev_tdata; // temp data - - // Define the blockSize - int blockSize = 128; - - // DONE: __global__ - __global__ void naiveParallelScanKernel(int n, int *odata, const int *idata, const int offset) + // TODO: __global__ + __global__ void kernelNaiveScan(int n, int offset, int* idata, int* odata) { // Using the Naive algorithm from GPU Gems 3, Section 39.2.1. - int idx = (blockIdx.x * blockDim.x) + threadIdx.x; - - if(idx >= n) - { - return; - } + int index = (blockIdx.x * blockDim.x) + threadIdx.x; - if (idx >= offset) - { - odata[idx] = idata[idx - offset] + idata[idx]; - } - else - { - odata[idx] = idata[idx]; - } - } + if (index >= n) return; - // Insert the identity and shift right - __global__ void makeExclusive(int n, int* odata, int* idata) { - int index = threadIdx.x + (blockIdx.x * blockDim.x); - if (index >= n) { - return; - } - if (index == 0) { - odata[index] = 0; + if (index >= offset) { + odata[index] = idata[index] + idata[index - offset]; } else { - odata[index] = idata[index - 1]; + 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) { - // DONE + void scan(int n, int* odata, const int* idata) { + int* bufferA, * bufferB; + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); - // Memory allocation for input data - cudaMalloc((void**)&dev_idata, n * sizeof(int)); - checkCUDAErrorFn("cudaMalloc dev_idata failed!"); - - // Memory allocation for output data - cudaMalloc((void**)&dev_odata, n * sizeof(int)); - checkCUDAErrorFn("cudaMalloc dev_odata failed!"); + cudaMalloc((void**)&bufferA, n * sizeof(int)); + checkCUDAError("cudaMalloc bufferA failed!"); + cudaMalloc((void**)&bufferB, n * sizeof(int)); + checkCUDAError("cudaMalloc bufferB failed!"); // Copy from host input data to device input data - cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); - checkCUDAErrorFn("cudaMemcpy to device failed!"); + cudaMemcpy(bufferA, idata, sizeof(int) * n, cudaMemcpyHostToDevice); - dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + timer().startGpuTimer(); - // ilog2ceil(x): computes the ceiling of log2(x) as an integer + // DONE + + // Calculate number of levels int numLevels = ilog2ceil(n); int offset; - timer().startGpuTimer(); - - for (int d = 1; d <= numLevels; ++d) - { + for(int d = 1; d <= numLevels; d++) { + offset = 1 << (d - 1); - naiveParallelScanKernel<<>> (n, offset, dev_odata, dev_idata); - std::swap(dev_idata, dev_odata); - } - // Insert the identity and shift to the right to make exclusive - makeExclusive<<>> (n, dev_odata, dev_idata); + kernelNaiveScan << > > (n, offset, bufferA, bufferB); + checkCUDAError("kernelNaiveScan failed!"); + + // Sync threads before swapping + cudaDeviceSynchronize(); + std::swap(bufferA, bufferB); + } timer().endGpuTimer(); - cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); - checkCUDAErrorFn("cudaMempcy to host failed!"); - - // Free the dev data - cudaFree(dev_idata); - cudaFree(dev_odata); - checkCUDAErrorFn("cudaFree failed!"); + // Insert identity and shift right to make exclusive + // Copy result from device back to host + odata[0] = 0; + cudaMemcpy(odata + 1, bufferA, sizeof(int) * n, cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy from bufferA to odata failed!"); + + // Free bufferA and bufferB memory + cudaFree(bufferA); + cudaFree(bufferB); } } -} +} \ No newline at end of file diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..85a9b61 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -18,11 +18,26 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO use `thrust::exclusive_scan` + // DONE 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()); + + // Copy from input pointer to host vector + thrust::host_vector host_in(idata, idata + n); + + // Copy from host vector to device vector + thrust::device_vector dev_input = host_in; + + // Allocate device output vector + thrust::device_vector dev_output(n); + + // Time the exclusive scan on the device + timer().startGpuTimer(); + thrust::exclusive_scan(dev_input.begin(), dev_input.end(), dev_output.begin()); timer().endGpuTimer(); + + // Copy the results back to the output array odata + thrust::copy(dev_output.begin(), dev_output.end(), odata); } } } From 3048bff1c7f5a4113a7d822075603117e510548f Mon Sep 17 00:00:00 2001 From: Nadine Adnane Date: Wed, 18 Sep 2024 23:57:38 -0400 Subject: [PATCH 3/3] WIP readme --- README.md | 85 +++++++++++++++++++++++-------------- images/graph1.png | Bin 0 -> 68203 bytes src/main.cpp | 2 +- stream_compaction/common.h | 2 +- 4 files changed, 56 insertions(+), 33 deletions(-) create mode 100644 images/graph1.png diff --git a/README.md b/README.md index 381fb9a..cf6f3e5 100644 --- a/README.md +++ b/README.md @@ -10,77 +10,100 @@ CUDA Stream Compaction * **Processor:** 12th Gen Intel(R) Core(TM) i9-12900H, 2500 Mhz, 14 Core(s), 20 Logical Processor(s) * **GPU:** NVIDIA GeForce RTX 3070 Ti Laptop GPU -### Results +Note: I used a late day on this assignment. I also need a few more hours to finish up my analysis. -Note: I am using a late day on this assignment. +### Introduction -## Part 1: CPU Scan & Stream Compaction +In this project, I set out to implement a few different versions of the Scan (Prefix Sum) algorithm and GPU stream compaction in CUDA. I first implemented the algorithms on the CPU as a basis for comparison and to reinforce my understanding of the algorithm. Then, I implemented the "naive" and "work-efficient" versions of the algorithm on the GPU using CUDA. Finally, I utilized some of my earlier implementations to implement GPU stream compaction. Through this project, I analyze the trade-offs between CPU and GPU performance and explore the benefits of parallel programming. -## Part 2: Naive GPU Scan Algorithm +### Features -## Part 3: Work-Efficient GPU Scan & Stream Compaction - # 3.1. Scan - # 3.2. Stream Compaction +1. **CPU Scan & Stream Compaction** + - **CPU Scan:** Implemented a straightforward exclusive prefix sum using a for loop for sequential processing. + - **Compaction Without Scan:** A basic method that filters out zero values without employing a scan operation. + - **Compaction With Scan:** An optimized approach that leverages the scan algorithm to enhance the efficiency of the compaction process. -## Part 4: Using Thrust's Implementation +2. **Naive GPU Scan** + - Developed a naive GPU scan algorithm following the method outlined in *GPU Gems 3*, Section 39.2.1. This implementation utilizes global memory and alternates between input/output arrays through multiple kernel invocations. -## Part 7: Write-up -# Comparing Implementations -# Performance Analysis +3. **Work-Efficient GPU Scan & Stream Compaction** + - **Work-Efficient Scan:** Implemented an optimized scan using a tree-based approach, as described in *GPU Gems 3*, Section 39.2.2, for better performance. + - **Stream Compaction with Scan:** Built upon the work-efficient scan by first mapping the input to a boolean array, scanning it, and then scattering the elements that satisfy the condition to achieve compaction. + - Efficiently handles arrays that are not sized to a power of two. + +4. **Thrust Library Integration** + - Integrated the Thrust library’s `exclusive_scan` function to perform stream compaction utilizing the GPU-accelerated primitives offered by Thrust. + +## Performance Analysis + +# Block Size Optimization + +# Scan Implementation Comparison + +# Stream Compaction Implementation Comparison + + # Test Program Output + +The following test output was generated by running the Scan and Stream compaction algorithms on: +- an array of size 2^21 +- an array of size 2^21 - 3 + +with a block size of 128. + ```**************** ** SCAN TESTS ** **************** - [ 3 7 33 47 27 3 16 21 20 14 40 30 20 ... 29 0 ] + [ 47 4 13 41 40 12 10 35 21 19 24 19 8 ... 46 0 ] ==== cpu scan, power-of-two ==== - elapsed time: 13.8246ms (std::chrono Measured) - [ 0 3 10 43 90 117 120 136 157 177 191 231 261 ... 102740616 102740645 ] + elapsed time: 6.6257ms (std::chrono Measured) + [ 0 47 51 64 105 145 157 167 202 223 242 266 285 ... 51377593 51377639 ] ==== cpu scan, non-power-of-two ==== - elapsed time: 13.7131ms (std::chrono Measured) - [ 0 3 10 43 90 117 120 136 157 177 191 231 261 ... 102740526 102740542 ] + elapsed time: 6.5007ms (std::chrono Measured) + [ 0 47 51 64 105 145 157 167 202 223 242 266 285 ... 51377490 51377518 ] passed ==== naive scan, power-of-two ==== - elapsed time: 4.41216ms (CUDA Measured) + elapsed time: 1.49398ms (CUDA Measured) passed ==== naive scan, non-power-of-two ==== - elapsed time: 4.15386ms (CUDA Measured) + elapsed time: 1.31354ms (CUDA Measured) passed ==== work-efficient scan, power-of-two ==== - elapsed time: 2.04285ms (CUDA Measured) + elapsed time: 1.48496ms (CUDA Measured) passed ==== work-efficient scan, non-power-of-two ==== - elapsed time: 1.68448ms (CUDA Measured) + elapsed time: 1.13254ms (CUDA Measured) passed ==== thrust scan, power-of-two ==== - elapsed time: 76.0596ms (CUDA Measured) + elapsed time: 0.857024ms (CUDA Measured) passed ==== thrust scan, non-power-of-two ==== - elapsed time: 46.633ms (CUDA Measured) + elapsed time: 0.697344ms (CUDA Measured) passed ***************************** ** STREAM COMPACTION TESTS ** ***************************** - [ 2 3 1 2 1 0 1 1 2 2 0 2 3 ... 3 0 ] + [ 2 1 3 2 1 2 1 2 3 1 0 3 1 ... 1 0 ] ==== cpu compact without scan, power-of-two ==== - elapsed time: 26.062ms (std::chrono Measured) - [ 2 3 1 2 1 1 1 2 2 2 3 1 1 ... 3 3 ] + elapsed time: 9.0749ms (std::chrono Measured) + [ 2 1 3 2 1 2 1 2 3 1 3 1 3 ... 3 1 ] passed ==== cpu compact without scan, non-power-of-two ==== - elapsed time: 26.0387ms (std::chrono Measured) - [ 2 3 1 2 1 1 1 2 2 2 3 1 1 ... 2 3 ] + elapsed time: 8.3302ms (std::chrono Measured) + [ 2 1 3 2 1 2 1 2 3 1 3 1 3 ... 3 3 ] passed ==== cpu compact with scan ==== - elapsed time: 71.3503ms (std::chrono Measured) - [ 2 3 1 2 1 1 1 2 2 2 3 1 1 ... 3 3 ] + elapsed time: 12.1135ms (std::chrono Measured) + [ 2 1 3 2 1 2 1 2 3 1 3 1 3 ... 3 1 ] passed ==== work-efficient compact, power-of-two ==== - elapsed time: 2.31478ms (CUDA Measured) + elapsed time: 1.50966ms (CUDA Measured) passed ==== work-efficient compact, non-power-of-two ==== - elapsed time: 2.08794ms (CUDA Measured) + elapsed time: 1.41312ms (CUDA Measured) passed``` \ No newline at end of file diff --git a/images/graph1.png b/images/graph1.png new file mode 100644 index 0000000000000000000000000000000000000000..f1f9204157d117b7aa0db9f925e4a4406b492847 GIT binary patch literal 68203 zcmeFacT`m8_C7u_8WYqQ1rP*6a6}?V9S~{K;k)0N!C-FgXRY5_zd!G+m37rQbIx0Kd7fwQ_grr!#S1&=4$xsR zm>m~orBpDOpNSaE7OkIuf`4(!J{$@E^MkF*g>#tf#)Bj9m(3<;FP+6;as#$48*G8U z|8hfC%NBzXc#i)20c)g9!eG{tFG`(NbHa`FtXW^3TV=Cb{pC-MP3)3mEV6-bc6&Z+Gyg&$^_D9T< z?fCucf95#gmbH$h+`*T!{*p4&}Tu<;p0tj~Aza;}wxR zZQ|l`DmpqkBm97ifrG=>;!ab?2IKPN|4$2byjg(zR z)E=6KKkMi%a~5PyJ4t=sCF=7?x?$d)X^&ho+Xe2*$K(u7r#^4ivvG?KaAdg7s94TH zJkerlDx;O+`bsU#+wfYZOY5RX=|#7WItE?0URb>CBFO3VleU{_4ZbTdUpU72Q?B%d z=>=bI-Ql?TmFp58>Mp*@=)_0LBrL1kB`D0@P3EGPj_j~}FC0BDdW|Xl=y$g{?Kb~Z zLQCg(8aI>SwmWUtRAEPI^=5bO-MjY@)AZZN=kKufc6nrxHdXIE#Lq6jmQ+~oW2w$4 zVPk0_5fJ$7r$U@$w5r`wOq)>AA$;WL9r9X(FYjz?DsPhMbGZq7xcF60Il=O{hP!u= zyc?@(ldqIk&KsfLIdKf;`3+SRRya6?BH#J`@*?6>famz(;OD8AV4Zt;Xg^)Lgc0>$UArP9?m70Q_1`cO6frCUKO<%q<# zES<+fyED8Q9&@#p8(@j~I_fE7Va{2c-a>}b{16A5Q<3?qg%P%A8RESLo6dJ%^bj^N zG&Jm)EGnu0G6KO>N|~cz^g;-uk5bsk2fh%zN}Rt{lQlVA!~bfj1Gnd9u)IWBWC>5W z@cX00J(lN-o6ajF5bRZgJNRvyz6`5tv0g2)8!65`ihSnSJ5{((d^ldEb5u5ScdNhV zSyuH5_7b06$b>I`yHT>wfx;i&X{TLuxs&{{g-B2+<-R#MG&Fr7KNaS3mX}!H&>%jy zdNjIQU}AoY{y`fv{tZ0E?(Zw8p}qCpONIntsNYMD;>+7Ed2Hd zh}sS+xF0>Gp?|+-U!;`PeH*I}Wm?Q^tbx~JeE0Eo6g|QUDk}2_cF%YB6-3cXUMO30 zx=c#8-xg1ZyQ3^|FRjz`xQ5<+QvWLMO%bJ*43O?V;0nQ-E@sNJyA)PjH2 z(&y?r4u1YCpWH=u2ir$$9f@a3>(_JG_CWdcJ^$Lj6mwc-%Ww*Nt&J0eW?lCud*kG? zj^=l(G>){U#xdoo^6QH|fGL_(%(@q?9%k4hX+qWvJX)R7Hgr7?ANh<&dROxBf=`oL zpy6d@|4wEmCix1Sd7t}h6FwhKd{Q=Ex}qkkdN$AAPnK6sk6*@Yj37B=Y;POWC)l+Z z(~#t6FwX;q)g@OafWe%>?@-lZT@`X^{^gyl!rf$6)BeH)dVHYMm@%seu>yB&{DWs@ zQXRcbJ(01%mN7?la7CEyQK#h#mOAmQ{^|)c-sDF0<5m+g3N5LPwKDydos;85xG=n- zeMm>{QE88_xNv;mTv7dmu5H4J)^U@JBXs9_Dm>h_hO###dq~!~i4^KAu=q%nzx{06 z+|=%-QX;BIQV(NaDsPrO%{D6zj=)EU5~f$(ff<5AzjNBr3(Y4kKP+4dUzKiY7(LGa zBJ4nhLZMj@kI;)=VmssA@q>jr#g7&GOaiW{Xrw!}Saus8%OY7x8wOP}DH~lsWS$sm zRyaNXRfm@-FfO-G+C#)uE6j9GqV9O(6_L*7x0cP1MU|g&EH3tV1l;L*z%`4zHjs?- zZdcTmT%h>*^%9r4TBAWsa2unI zPMJnbsqy5TlXcoJQ%;^e7XB^TM&7xl`%R9eZECFzW}9{;UJQ-sb4%L(a&kNs*1FFv z|AG39?ZKML*@ZUf zb7yQ@fB$U%bZ4z`MU6`5q=%cLYFwUdX>X~=@Zia6d$td*-AUQ|r7KeJj+Zh*F5zSc zyo?Fm6MQZ!_^dSRW)VW<#n&A})t{WEH#;jt6p)OMvX=C%4i1U5ocee$!1otI$iTor zVSkk?*HKn!)nH5h+C%+C^pfYz3Ia}gtZ;3vGTWAw<2+a&x%{C|X5gUA$CO!T8woLS z5sz%O$BAPO8b{=gW%cJTIhVG7cC(+>@@=|K)*P4Hb^C_f68mhoFR9L~=zvV~FH3E+ z-AUcy*7zcWqEuw zZrsX_W0z&MFuvzts#TSv+Gt+eh#;RuZ#B$PhmwY$WI{+h!N%Gp=3Sf9Rw-ev7+Bk< z=6qb`llY%)$1R&Ljz-*<&tly*l&+^TA+kE_sg-#1*Z%2ui)9I0?zEY4=8S@V29*A&IMD2I9X2QG@=pH^4ntrJLV(TJS7PXYKcNLON1=W z$-uCO-XM+tD4x})U++U;oT+d&N?u(h7f$;kbEf5QJWoD+e6)G zz2j6foI`X!V%0Uh_h_E-juV<)#OGW}3za^#vK?Ppqf&VxuxrzkxP+^W6P&Z}8oc`t z#oe2?8cDXQoFK|t{LC7%bymDVTlm`l}&=x01$$iuFKk6rND5=4NC`*#TG9`?f(ZZ;UR?1kC zyX>W_4uXmb&m9F7&4hv#*EUJ8u>Po1Gt&9V+?@RwD_bc)-4GwEiB257cbsL0^x1?3 zaXVMsDO2}Y+2&1hqfVA;qx+NkNxaG)7gzjxj%UXzA11#xCJLmPYenu!h)RC#GglO- zcJc9>4za90!A|X`qiPK7i{g9xlM0DGa}}WpXF{}Jq^kBg3N@ZTm}tcyuGk~zZsn8x zRLJk2c5OF}Qz5b1sS3!J$H}r~(Wmx=oIW=?du?|`r^>(`D~*7k6J35i7ORmkd{1p8 zfpl@)KPmgiSfA}X($u{STRe{4zUj7<91aP|L+b2ty2JS7M)PKZLI|%4r;FXvg8ZIe zg-ae4xUN_|I4|OJm6+fZ$Pra-*%Kn*a^!=f8&0LhDzU1+KW|qPZ{n5C<&ZWL)(mrdzYM&0%dlNx97{4;M4- zwO<^0aF#r}L%MM;{$Q`i#ZY|Tv#QxbT|=RVPPo(Av&C7(HwWZa<(PbJbmDYo9L4Op zNn;)7(%Kah+8zeEtjzmag(`Lu+QKeOPFuFuiQuL7^GIyN`v^{IaS=`_qA{y_wjnr=(Ed-f#N z%PjF#xtIJ~Wdr-mA_)qk+8yKekB?+;$-4HsDH6s_sJS61ToRT{cQ?>Lb3eOBHdaS^ z^7yRaiPSffpWB5hwMI*JciwiWP7nFlh3@Yr+r8^c$yC!yA$k5GNjLQLh=%Y*%YAJ; zlU=xdVkPWP?Klhes4gooT?}>h@sQa1asI13gw`GW%FLDLv-+h^g>DsA4ZA2_bbVKP ztlHUQK5XpK&JzGUwD;DWE;y%H)LB2yASr*awIeQBG}a>(Z5qK1Hp!K>ColjGr##Lp?zrP3XUty73*j{ywV zgWGpNr6#$@j4Q3P*Fr1RuT2Om;hGh<@Ou0emxLaZWpjS>gf{H;4x0x$HNJBi4mySZ zwrsCXo(*5zZYRbr0da>#L(p4Ad6_G1QpNk18B&q;&@sEz@S&9ZZRCcAhJm^4aBc4t zQoB$g8I~@GC~!pG%r{?5@`$~})-0}8t9p+zkNxA>_{sg&<>D>lCZ0-F6Z&Quk7eDy zz8h&5lFZjBNR6I2jXc-Okc!#lE~Pk#HMhSV|J03T;8vlx%UwKwa0tiZQGxT}*Xgw_ z(pm9)DLQtXRrZ?UX^-r-uB`6s3B|!}#Un$I6gSo1*|u-WwbOU5d9e$paM*eP*2*IU zpHTnyl(A7_oaF4n{(OQ$WGG&EX%U}a8dgommKJwR;K@K%LZ69ZsJh`&rh|)p!_G9z z5+4to&W?Txpm(6iR2-bFZrW@iC9%V1!wyrKSAwn5Y>xI|#S|H~X){aS4_;_cN#2T& zG&E%S^Jw)0e~q;+S+kc#BwJNAvIEz6qn^2eKFAQ2sg8^85C@s@kq_T-r0PbpCuL7a z6St;~-w-a{p_0D6`SWy)%Zk582aYtXy_dARdZ>JM6&oe>(IO%_TXfR4pf$tX!|ih6 z1Nxeia)v#OuO|#)`Nk;%2r$U{*L4qqPd&0B{q9>kiYp&dK6$UlN1dh57QOSM2o~NMIOyfP?Y1%UQm{hJ4AHi?K$w>Ks3zGr9tG#pjv$3V&`P8 z!0X?Eb7rYMKC(NQjIS;{lGQ)sx+kknWwx=U z*}5T4`QdXC-L5R@%)4!5>-un>89YT{^#fIN2lR#_pcKMYNStH?7LqUFG~?tVw{HSV?DTBstppVJwnypz4zq`DA&-ZxhGub+7XX9Xw@sy~lm! z!>u&+93Xx(8CX)o6C(G?(Iuz`X(`&^Nh2wPL_MLNHo}U~# z%;f!2Mlvb8I}n$7YR|Q#*H$Kioy+rF$v)P(9QDTETfDrM@wFbm;3(`wXObyz)t`9s z(zt`6r&{AjW8G4R@#4-g2Ocen2LpH)bkt9^FhrG81U_p@SkDQ(YYXf4^J z6d5Y_(&11qRB?9@Gxg4rOTx=p{Miw5M&(4|KR@Wmi@zs%Ue_yCz@4T5v)f zPfK7?)C5*1z(jV4X*2Rz=E~KU-8K2FL**RsU&*9G*^xoL`GA+Aetk0Jex{6chGwgy z&n%7{FG!cEZ1j{%&kW|U4C%RL3KB&r;bP?jlbkvM<-Hem((sluaYu06Z6%)bmT*?jNwbZ!g63M3SVA7 zTX`6z^x1d-`IybJYVccWwhM{#&5Q39=^NS8E`A_zx>x|E9_$Qg(kJ&4SFe(6+;n;! zk4;J}PKt#BJDBi@sXXf4TRtyJ3O#aE`U^jQ2HJUB6e}ca_I;zXpPFvbg z&~Q`U?5tsF3<2ozYa^FYXa2N<(sw(48ICyVPa<6ka89-JXot&)jSw3tzsJ%C?SnWtsbX<9`T4bWoWbE!5 zc8|FJE~&nWJQF<6ct*MSuYV&NGj*{^9PcaqxjRVc02*+KIuOf>tfYy=L+*PBWAvC=1j|q zh*EVhPrI9Sx?)m0ZAhr4!QonL)n<>;0+&GSg%GXwqlZZ&2h=_Z9$ z+*e_pxIkV+bB@1pRNOk=r ztI%NV@ma#Nw{HUo3Rll3JZbezLTkSoA7l^Hr&2VCd;^@2f=cg zF1uOo;JiZ(3`tKCI$N-{c6pmra+&1Ck0EUF*5FSgHAM7CGEwR}6RO7Cik;+FL+o;h?bGrtwDI_j`x3_a%oh<}92JiQN9VV4f@Gfk6&YY(=7C z(zdEA_m;}C2;GU5-Sp(NYZ@n39z|HYms&$5z}BkEDn*bAqE%e0drd+$|vhc|#nJHhJQ@_^H)~@^Y;K`@iPbWdYU&bw_*>6iZxzk1h zT@>F*DLuy5R{QSj3>EI1nCS({`H&;}uJBKk889cD=f8X`X}P;caj}ViEBpgZ^*iI+ zp{{6bY&_*Ue6OS34VA%M(a+-NsJE(0 zogU@7-#@~r$aDv(>O%E7_l#s8>K)SfNh90!zo2c2l7g zs&VX|th8^Q|D_!EII9bnR{S%422sj+aI(K<;TT~IZZ$g|5yk8{J=j1gtD)297pm{8x>XXr68t@!A4@#oz22i=y{x|;+QwCqqlh6er%xZE zYA@x(cQOPod&+M3t@b2~kDzo_zoC4}o3Swq4^G+z&;L>i`zX4TS>kGuTs@SirztLs zzPuYv*xm%PQ=PW*Ak@JZvS`1+yt=G)cPK`}-6f057ST2h+816QQ@LK*va+#t)(JqR zz217Id|_e10!o3=j~-?3GwrL)-QCdAqO$t)`VU?t!w()D6eVbcc?&4q52}YH>9o2$ z0A=8AnET4)`XkX?hy48gv-ap4wnwC;rP(J(_}_%EWMN~&&-@GRGX1Am;+L4gDdOYv z_c~S&vP5--nOa$8SF|X6f^s)mBf(oiLBS{^BV*zE?U8Uj53Qw4-|zcbWOJ+BmY$j~ zbJYoCn)1Anf`wJ$tHMyHDVha+tyKY7II67Ww4`w6Wsj z+}zyl@!Tz4c@mFMy95wx-U?SBL?}mzaMgp_U_7212;CYR z#_}tE|GTKjNUbYgkr5GZpbi|_yRR!@ zKzi~uGtAw2D8pho?^-vJxh!g>cx}EZXW))^IH9O903UAc9w?s0J-806^aVrk*Uhc0 zc%oxsxIzJ`eQNsYRI!BTF_OFCM9D3M(lN%~(Ae@*p_25Jz>V>;Vkf@sn3=Krx=79m zuLGDhZSv!Il^Nn7=?tyxON!%7pT&#g)-%HytvZ2x)H`TfU4eE!*P{~5+`%5DBeIo9^C=YK5kjc3xeu51&o)LzOP zkDy&thcG9$R!WOiPVmeePw!$v+HlK~f}@p?*|MOuhA%lF5#R zT%+3LE6xUe$;+X%tA?c@yH{6N7ZB=!A5QD!SvWh3!n~URyW}B`zIL$sKz#YE4&REw zwq)F(sU^Rs5bd|Z2+O{@w-ez)o~v(w&5e@yg{;gKEjsO>)5F5ll-*PKEtk60i2di& z&!Q_2(0<3H9Cn0V!_PFt-~qEFj<@Q>J|h*>PgSn=|8(UapR@|Zg;#y!*U{svkN$KE1Kf}tv%mBh?0=86- zgsBY7oz5xHv|XGTS$_0lwzPMJ$8}~n@J-mKE1nCy$G7oB*D+_UE%mN`RGcc1{;pTr zB{IUA*`ZG0c=ae>u8YDa1O&6jb)VfZ&^0iISpp1PnIiI<$JqyZ3=PbE&EhPak*KVb zSR@YxM6Q%h2S|;>KWq3 zyYb;ik49gtFWM=F-|sRAEZZC(36OhP5nDnh;d8umpnh8$3-}?mx2R&+S6A5 z0I;L|@MLQ+XxGV}D?!y^9)nra*ScHZ8LslsE<$u8FE6kA(L8x5jS_ zz34G|oUG|FyKKo#;w>OHCb|-)J0)v-Y3%Dv0Xv|3*q^LW9N_YYy(PKOuv;)Bm3_51 zsQ1ScVy^ZaTGub8;O<@wF)FFAP+cP5S@$XqRd^~?xX3=dBGw=3Crl^$s(sA^0&ZOj zG4aYLPm2(G*uJuEGyHU?zKeoMr>ZRTP6 zB3|g1r1z`<;pW}7qsp0D&^3JFhf@nyAD70)18VYOFU%d6=N5bIh_{YqCW6}@^?W@M zH{Dm7bDUXE*g8fmLOg3BmbNhI0K5x4U~6l;koN}xJV}T5<%!hT(Du-= zvGz2VNLKY*?7@u3UejXOlTG;j!VDzdO9~3diTn-TS>a(}VN)hQ=nQ-aUYcz;P+^H6 zI4|73^7kq#r*NG=q6`2&vbHfsH29LDVh19Ci?qi=xHM6LDu}5R!hyz970WGNweHGCq$6lx=7ph;!TIl4H^@z4?GIt2u zpWir?frK6l3nVH$2j?uW> zb{_nGFFIFko%#=}PN!ia(fe3&b!b%IW&*MSrcItl+2JBug*p-Cmoy7xltSNB3(Muk zs-(mCa?4dMrfF=34=5r=a|kx98|*z8cFfxbrLAw1`wD&1zmH61@s4s-Y1X(u?vVUm{W{}0@AB^11*UK7M}oe5uLl)6^AZ$~MmEv^{)Njg z>)VIgNR-~xjmkNUic~f__)W-#>DL{9eG>NU+X9l&HFf1|bL5OyXb{s=TXT#Z55I3% zpZu^SkkOKx3myBqKhwLSX^1*3#AX6Q2tWcPeBMgJBG_@$@hIC}+ym?5Bck_V;oB49 zC1Dx5^1I|lsvY$%!T`RrqD3gfdOZ^T?^ZOt10L@(`nMJ3YBxA;>m6F(J&Q-MO74@h z??u*!w|{L+Q^Qq*_-z7|E9kd`L-(xg-Ad%>Uz1x2pZ6wD4r2US=a`#^ealxaEtW- z{Ot-sbcJ-t)r7Gi4zP3*b3mv3N$LV(AkAZ*S{gDXcPX+XAME3xaFq zUz?+AcPn(4Ex*XsKpgr3;s;KA`n^pzFbwA)Ya?zK`laWS=W} z-y+fm>D|{5?pRT`zKQ8)V2>U%=&+a1f99k-!sfHSdj=bIl=?Gh2yKnH*t&04em@sR zz3xn2%1Lnhw@V}8aWp%p&uDim`r^L627G2{mGFC$BTT|D>6tFTyS}iQ zp}R@gY6rQrO$yyrUlk$v1+Ba4_8-9tIY$1gG4Z{DF+R!vuf{})cWNJyoqw(AuZ0JU z!D=C(v2N##CtwvG`UVkwCn|zaNM{=QwPW4TpXwA4q5nO%SKPPm2s7i)@Vh9T+)|95K|M*X$JfeXZIKaAQ0$=JnZT@_=>_=Xo0BJVikNMkfMbX8@fc7dc37;S$ExbD~^-q4wA;Tzm&S+UAGLYDl=^#gu;2_ud&^O-)UWR!pjndlR zv6FEp)hn?OS22Z}ncL~V<+rLj8t*v6kAAp=V&{^FFYo=&?Dkv+tq_xREn=NP|IBUI zFQa=?$2AV*5-jba#f070QMGkbCa)O7w+7o;PT4Y4y~WroAmv}k#&ax1PR~QqhR=7P?Dbn>@njxZsCE4xmgMyU(hJ52 zWD#X5%#8LPncEnVwgXea#D9k6V*+%7vPbFVv4)DWY4vaO2mKwRA3L$Pq6r5pi&pr< zNk=}EvDUMw{t`OS(WC|2qPQ~H+5KQr7cyaF6(dS1HB{8;HJ$<^$+w39x;V-QV5Pgt zOlSj~^BPARVQ04i&iw-~d1WQT6%4z#gkKBqy@(tW@Y3ui+EQ7duoNHijS`@!lxaOn zgS-xB zfW-4OZXJ-Ekwm{2S9aCNuucu(cplP#xF{s7)9c&ig2K*l+IM+);bU{ax=TKF217lg z7gM=>V>)lnN#FeT)@6+wVWXXsx669SBw;ZkgI8Q-MJKM~Uz5r`^Y+`BfWVA!B_ zpmW9sp8hxSVO{o2EY7j4uY=w0CAJkw2 zUh+CYxk5LtGH@ec`r67XxejZeBi#7>jOPDoP0)W2*B})0_1y()RtTbin>I!#G}j;O z|GPEC*P$!^*s~=ZTapC%BX4^6kqeDqdYPevqCep*k{NV3k`V#YIbE%Bnd*~j!LSH? zC%%JpCugJeOaMHrFO(WkPz;}%$Kt@ydUo?}H`~D6;?Yqwug5e0UDg@ z=Tuhg#dKNyAGVcvM(b`iba!WQ)*-wpoA7BjJ{;kT1UGt?gVeRT8%Nl#tA)cgLV+75 z$ap>1U&Vsr43{h>7)Yn72!-!GNIrFKZEd#FApBr{YTppR==(QpqL)D;*LhRWUaoT1 zKyNGLBd^DkBSmqYRG4aPxLCe{?md616!tg-)E#>gFVr+SV32D6Oyixlz3Y*eZ+i$H zJMgJC*t|MIKAZ;+0xQ_m_GF~%?CW)BK>u z^6SIe+S-D1^&muhWWX*a4geb`Py`-6yx?kg?#6k?!@Lub78z>5%E~&W;|UPNOBr}2 z&9VWGf3xftkGWpm1wpFJYAJc%6;VGK^?x;U3t-ppwBbN{&hb|4OLUgFa%d;oQ#=?@ z_jiFb)nl5xNYjojn32Z)=n;=E~mdtHKimA-CD;CdC+DBss-5`*xXW^B2u2-j^!cMN^jlSr&1*852LPa@H z&GnSMF}Of?mL|#sCHq*yv85F81jzq7boTlHb@D^-z^<5$z{1CpCZG+&cfxoN4-eh5 zKVjy}nGdf-jEfaWB7a~U+K>$_pbj35&==Fv45^7sh0=n+K{sd|@l4Ryg6M!93T@>p zlTPd7MQN0l7`+G~_$~$*5cV^q+p!I(dMs7Xqmjn64(g_~}GvFi#2MF7Dt3IGtmgK|tkpx>L`^I^&;&!NN*noP7EbrX8|6uXgqEdZ7xAqa8&4MNg(M25WR zV62(@bX90-SGgBe>qW^3o-kH=XM#+H9Mb@6I*MGVS{tJgvV=WFkZNL~nwR&VY)$_+ zWC=?9GG3OUpa2mFUR67@QG1@V>$`ElJ)3L$p`h))yGz}N6p2BT2; z4*;w0ik|Vd1N1Z6AoD<$TwDi6B-D55^bf}KQ$yBWUXE*-(l=n z=8W(g)cr0*p#+l*h;x+6`Vndy1NOD987!igtaxSjfVKYzS$lgKEuw|?ztVn)uRmhW zTt>(hiYMPh?G1Qg<^g)PxIO5TUqSg=m-@>8Y~|93U6^XTf>caGxTKSPPk+CC;!|6| zo=xIZ#o>2k#cyq?@=}y9SC27L*dS^xK2xspRyR=zFuycCO`Rj+{+Qut(0NTLYMmD1 zIiBp!uBlbznkDKrKckt<^rZHBJfcJ{;}J`b&}ZXw8TC~i3e-y|bM*#!1(nL7%+xe} zp{aN8ok$eUE14+Q#HXW2b*TUyH3FHS&K4hL*TPwC7Y&mTLL9r7%EOhUDK{&|X2yW= zku%DZ-&ovIzvtt$ z_fz+{LtG+Ze#?6(sM3oAMRWbXlIpn26+>=Pe&#hw`M3CG&^@4fnKO-QJe*XQQb(EV z`57>G79w{!WpCj(6mXCE(FD2hEn>>l1@hWq&MwN$Fztz4ty|^g<-1WInYehYZkZE5 zv=w9yL7!tc0?=W(xw!>PJXfD7c`g^2cM_oR?RFgFW%ka|Y^}r)9^b^g08^qqAEuOY z;jn~TL_>YOQB44|-j(|!!WSMKOxt(es<$FGD?58IXA?$g2Lxsod7kpRoCnKWIKhhG zb4D7D44=!y&~w)0XEU}2#6YiLsq?h)Y8|}mWgLQ3ZIFD^)CJFqiuP7ik;tL8P68=W z*k`y0Rb`FRl11iMh+83#8R20~42=$;+DM@l+E6At~~f;hu1rv&p%8S{yP*cXKx7 zTm7&F!&e9KNG?n#C@tJzH%sI#RGZ90Ruf9K8j#;rl|t}_(aD+v_}`?IhoEP&8#+O! zbRcVVqO5_$$1UMFHr&v?3zCC55Ch=jcL1ITC}Uc3M_7A%`!aCD#4M;h`6FHThH1Fx z=Kx1m>xmEu)P$U&idsFY>dDB0JZXC<5>rr3kK}B3nUp{THcucHRlK28DuB69*+}19 zU1VFxpqfhJIb{8h?ZNt8oKr-z45gzwC0qox?y?7IHc*N^)P500Xd-~2m0?RmdS0f9 zDi?aUP`}_?4#5|q-dJI1=?tTFSH3mR&4Od4s>hK+=&nO8Sm-}vC(fJ=!SU#pgn(gw zDmra#Z!aVUFXKV|Fp}ZW>)Tm_n)N9W8kJY?kAxt|x}K@--XkdNEZD10>Fe+Ct${w) zo4RLHgtq>A@Xg@h;FJkEfI!jGkc9GEE4y+&%9-FN=%e3-CC(_B|B#}qm7=Pg5{Rv* z7|kY3d^loqC@yK;XAc4mYbucP122OFuGV9@bPDFpECoASfnLSAzK25zWYn5s?L#^+ zSK|(mvIgaLRR7waBHLHYPZh)YPKV%KKFy4@nr{i$6lE>cmh+yH@3&F*(bCj>GqG?R z+i>ErAoPp&(qYb2R&nsB2`B_sjh>KPnX;AGO?f)sFuDW0vkDjmyr5*cV>aPTtJ77Z zM0IGr5Xn6NtG_~#Adnz-P9YVQB$kliBTM=Db9C%*0wH5w?jyw!kB5Og)DOZ)99D^kVI8);)qgtJS_|s?4o@rP_Mn%e-o3fI=oQl5`wyP4-LEiKx)EOh+rxm#2%n%kV- z*Vij}Tz&@!raKOoqONc{cykeo%{lHf#sshzgd0!=tC_k3XVDJ$A5|#ba}9bb2+;8$ z^}_fVifFDDnO7mkcK8s8xx|U%L_X==_-Jeo})3|A?(+SFtMONKh^J3srP z6Ca)((}txf*wzz=Nmj_X3mYMq8zG3hjX#!vU;|*SX_^&*HraYi`R0BmCY>uj7g;po zy=zjW_*0y>M7;r%nsx=6pa@2rTcM(m4r@q}hp8s;G!{(YIy9JxLA!hjLEg5nLenxO zwqwtA+Q}Bi%?R#LO()rZ^}GTulfoKcmscy&Th@EVx~;Pn>-T>ry9pN68m94}8{IH(p; zWmgB`EDvD-S_7K}DkovO{zRfADq8I|gl-UFN=b@pV_kG8s64%xo7_T)*;Nd+sKiR0 z#FZ&;soO(j-rkS zib_L`Ho%q@c4lU0XX_{%gSk|0#b8$d0(?xl6eXoNj|EYhaSw;VweWJZoIQ9-&l0?8 zv}p+;%nfmbLS(5n_;v?A@-U^vjf-A~?+46tnS?LM!i|oks3u&eze7WNV&E|X{QS9v zB?_4*AZ;>vbB9tC{Ua%|NVZ>JGiX~D7Z4LG)d(ww*Ce85{(N9>L?rqnb1afP+_b^{ zq4m(k2;~UWQ^;it=c^Wbf561wM;d*Yfsavd+R7{~4JrA-w77YAj1e4oiCldT#5XL$ z??Gt6pmT5x$D=48aE;U7aofw^Q7U69)dBo($=lq71;CV6e3``_e`h1LsZj zd<1sERVJt_4eE|4bHEBb#g$MkrURZD2uNy3jt4+nEu3E=&t1BF*%5lj(#gQ_W0C1n z8dOoujW!5)MPDHH^$9tF+?Vy}I1P{ki$+TQ3_4)I^YEwvgwmZrjghde@JdMCqJ(-% z6arTkN%*q6I$Z6$jXk54!iEB-j{r)t)?nrze7Hjm&6D9WS)`NDtSVg|BJQt;PO0${ zt2|*j^`M|gVe}A~r=Vy~k(;PplxlF|GysmhWM=Wub2fngN?rn?A594wh7f$B`Dnq(t#OaQbej)mrj0?EELxle;0(w8KjdkHImTZg3kiK5{;%9gz<@L_XF7 z&!c@!;1+O>39$UgW9%DjJvzA!s7OzbomCwYKkDqG?h|@7V(T?JL0G|ipngy)7O+LA zd4o1xnP`~9K?%2KxQ8PbVCn8}^i2?6$;ni|V8nyHQ^-I(s`n!kM?w*`zD3yw{Tc9f z)a-(}{RwF?8GHKj&f*60bme_y2xml!Viw#KuEKJ29Gi5f{YF-4 zXdm_WhEQ7DFiJ;IhFZ^>D_@3dX#2&eh=REiKWcxa5JyV9q>x7!y&J* zAWKEMh`b18WB9n>7Exe+SCrs&kI~+J_2Z->GCmAv&R_(+hN`?H`4y{hl63V=)~+p& zp_9@DR`w8LROm*~VPUiZ%xwn>v3u{nedOv0HPhrk%~_Vi^duqa$8c(i-P3_zS;qTj zJs`_CowF&T6rSc(_!99PV(R&Ozq0h!K!Yv@>vaF0SBK%^Ypf zK5QDU`3Gwq@A&Sy)JL6BPGcH-RS}$VV@{JYO<~R;gZbe2#T35^!Uf~wq!Fw8!iTjws&n&Y=2J?!3%pP>cF=Tu;U%>%1NGJf9jFeMW$-xV-EMRs zAYeQSZ6-2wO63Kp1gcz9`WUHc5mhiA(UZ7^!tux~ds}hj!b`x9!u~6sE)0jmi{9Mf zhGI`-d@Q!Lml5j0dQb*T{DibWBp|}NQMeh^>7TkoDa#vLdUHbYsJj(S;0cn@{v9mq zli7(+sRgaNgC}pqzeNMuo`Wrw8evg~DPu`eBUC3xFNo6h<&FZkc$$qB6cdZ?-3GVz zwG353sS;XcMoFirKC1!e*Ft3YQ-tX0>1U_FTiT~kS#cdJAUzm5glQ3=>p-1ANDxQe zLD8E@6$$kKjaIRcvOYoHHAWgU75^c?NCSu~sa)sT0iw<($q7Otu0tg5!mouJN;*2)^3~@&HlUWv0cNk!o3To&W_v~^+ji&pwFn^UZUr*fFIA} zK~su7^v2mSU* zT&PmhQn@_o>3EzH!Pkga2uww5#y7A2U03Q}#S+h00ajrGoX=T*_1TRh&|6=IDt zg{uN32KxNv*Vk>?w`MJLtg(DeeB(?ktf@13E7*`hM7m4m9*kG|BCIb(^#zV6)jUWp zwx-e3ct;@KO8FV+tm6JpDeY}-YQmzTT1^n9(9um)CL=?qAs{Ik z*SqAZ9AWxCj)V9o@D7*mB65kX$}z>DqpdWhh1zL*0!z6*`>%ZA zj3_{cn;&1^DImy}-jWX_DBMqwOLU@N1F+UG(XFkFF~jkvs3NR&oyZFR_W=J+0hM9^ zp;S4@D>HiEB!VY;kX@vS0w2+i2DmTJ+jx+_-XBj+__@_dGVN-aQ*(zrB-HO??%&Mt z6au8L<@j>dmBj$jGz1(9vGh;Kr=MV{6P#Q%BtFq}hsYvKrcn;>c2i1Wo2t8XMF2dO znoL$+*AbqEBjOlGDHqUG9@MIcgc};XLhp(Gqst);eu54Fx#J=U=p}2Q8T6142hk`- zZOh+S$TT$iPX&4zSL&>?b+1jW?CfSs{}zk^BKLCSnWKT+qppI}t}kT?w_89tn z0_fj#;a2mk$m(ZFS>RdGk_+U#1&Bs)I{OP4OfbC2zyh7K7C>*{mjICb78>ZOeRUuX z(6>$u9D6Z5GVoPCEB6qkND9L?-X>*Z?BY`LD46YXhde)GU5l3JFc5kxz;zG}k0C)B z+Sb_+s)P19${~^+GBwGSqTyU6GaPFRLszCLvb)zuRH=j}57cpwI0|avrTPKmj;)Q$ z1xB@a+sLa+(~lBOs6zevFR~~tvl}!NO6~8(7joxbg$C!vQqLIb0krLux4!E?k^YFt z91{Jcnpx>Y_~=vy!f->4l1?GtT{m3)w=##?DBO-SDoU=rN=9eR7IZu(3|4*6;_;Os zA;&9F>EC#`{Ys$r%t;0-#r_T!CBVt?O^@v;#Rkqqx4TsseRkWFv#tABaRj?9pIo-<}S4}&<; z{ZuP%tTm2xxu#CXQMmu>V1mN7#yCmO$JjjMhN$I4xa%|LFL&;NP<3kIeWzqx`Sc6x zY3>)-4@!6lER7A0*BZ*LEnn-q=G(Nr)tIMvqykq=Vezj+2@3KZ*}5c7V%2U1xXt|H zi;v94;Z~X;L~gGeV7g**&vPZuP5!A{T?8(RY)ES61+*ovV+hxYg6E4&?&TStTovjebZrT}`Z9(VptRy}M|scV%uT zn`Gw+!KTkmgM(`!Ym_r}PXQ5V-~95|i*5Kwcn^#C8bDvZY{QGmn%eB@_tS!N#usEN zSgzOP&MBJcJ@Ygn+-Nz0swkK0m2An^T!ZvldHIhiEOQ+RQ}w}zs_1ndzhdQb+;w({ z&4om9x=_%OgMjyp6_TvTq4>`wk2-H0pIol?#zF{ct_TXQU0k54XrcJJS}pn>qAKKQ)0LtmA*#rV4dq_FTA zJer?WMYnXjIc-8X#@5?dwE1>{sE#~?reSfCGNZ9ps)u*4h=8}A1S>nZO9VQVi_6BAcSb4>_S)AIcC9sK6u3N3 z?~2(MIPF~v943fPpw&+roFzdk;&D0%7Uyyrp>+MeQ7C9fUvG%25>s!8wGT{O) z!pqE;ZdaMc^iLLc)?%5)@$1`YtA*r(cv*cQ`A1Oz%OY3=S3G zv=7m6thTN=o^;gzrvnpOcWkihKq=eyKIR+R?0Y-c}tgmhY8K%eJuNj_hu z8GkJFn@;V>46jpqX~?wYpi3(L9rDxyz!e>UYU^dByo6Ee8%^Nl3~q}EF)#Xlzk9w3 z>=>I`PaogqPR*7w%8LlG#H`_%!=apURTx44P;G-1n8g138~AiSRfzcw%Q9Fl!Llt%Zp3Eu4;-5RskF$8IiL;jMXyfAa69P zEgci26G`IJsT%DUBt`L#^r@S#>BPscpIhLqC3hQs@B+aDB=E{kNkq5)YWk}Qd2Bn# zcjeOF2K&oYU1U_kwS!LsK9#Zt(UcUH+2YS=yH@CxE>%2o!7av<9ysl#=F816azdta z3Ahh0`Vn1kB%tT@cL=NdKg7&|O6M4y-jw@sw9npr2WGt^814%V&zcwyb#pyAkl|vd!67kXb8~ZDwtjsdf&U9_l>yJAeVQPm znE;)+gmS33n%C?kxMVo#{qX3S41HdeWN~~pA;|@IH(7=-;0NuQk^Nun`3bnLw#lC+ zfDRA4N^c^c>H=<5VjC9tyg<*zoFM(kDtE??1-znef>LsW=}9ZN+s*M~->-cC7g)Nm zuVDENE`p$3IhI(*I`5H70wl3fsmQWXtyqIeyef_(?wrXZ7xO;mSr5sBOiZ|piN7vb z1TJcde?DA%hh1RtdwqM{Yi`lJfY0L(9n(-xv<0pP`iEZn{IRayRh!x)4KKUqLTOVW zSL};OSOeVm1H~EHR8u|oJ+ZxQ@H(8J;=%=)&i&0Yg0+6mNOW2qO$^KtH1D6g7d4N+ zshfaJ7T4Qv!f~42waSidQ6Fo))aT?BpwwBL=U92x(BnrkFSn zXA8HLb49xp&?YZ#tULxz9(R=Lb?m9~RYP9s!cqB~kgo71%AfQjQC!aGP$6ii; z>g;=|v0p9C1AGdIrl3{^#UImJRhJA*RF`Hn$E+};6Zy?CF&zz@VlkhMX#TG*Qyux8 z5<46}h$i-W+b`BH?OfDDnnk@4qPG5D$7Y5e7wIeKgL($XNVs7}b5v@O+5nef>M>U1 z#HYqZQGIU1N0+d|INwZvjA^tSskMbdMSKOr}K_+>2JfVhbT8OXq!q*Rh!2uGXo+(E833|5WK` zz$KvM*5aM6gbO$|Zz86t`Hg1Kk-EUUj*6XaPWHflAxi4{S)*25L9YmM<-y&^CW(oj zzsoP~r&-b#d5HT9w8l?v!Qym&B|}Wto*r-?Cq%)(u*! z{|YXgMcM*Q^jNJScbil(z!qT1z~nGsE7+OA3H>gxRRJA(+l$`tvCBAWUvpWdA?&;c-q5%;h z&^Dwqzh9>la|fPkYQERoeUPh%Fo_M$I(uYpf;l|YuSizx1kh~E8YUGErm*ykicDH+ zOKa9xT4)1a(5L9elEs{~;bRg#=b4yqxWK8bvoE~^e<;`eNU$|+yDd%_Gl=(4NA5s8 zvB21-DB$pA;@a1=l?CuKK#Q0D5;?Nutpy4MDNx#ruovN=}lZf?cZmN=r#cxEYa zGQ8_PE=dCZ3UN>-GdvS5-1@wZxv~xJo;xHl@$&!*kkUrec7~*QxD;tiK73$x;^5P( zp=czqws^k>i#^|Lnm2Z)-7RSF&cp;=b9DVDaiLo}V80eBKmKv(ip_6dME-8p1zgLR zb+Wq2$#O3yySzqv2VAuzQB(*DX zR@yLXZ0El33A;>3{>UX*xkT0YsSe1mT4AU@bC24@GL#Lud&n<9NCr6FjF^Ya(6n>` z(fRRo4G8vmaUPRBL=Ts0?L4rwURiG3#rNv>V-*YpNAgK|4c8$<;7|?B7`gw*+t)?X zTWXhEWkF8vQ}i{^4}X0QR4y9yG12qDXW*gpDEi}_qF{#q!?H;3j*p{7_9n{#WAjNt z@TsCAxt)@d5g+dfNgD9NtNrQwKpHY2ugKSyp_CqN zvKoN+Mz~cSY9Jmyznw#hqOMVT(p^7;?7^pvB)2N#8Gnu=|YW z?Y%&Yns_#_Dl-7;Whj&s5FO-RvBJlLXZ&ppEaz_zH~ZjeafXmuyobwt)@RU~cCI4R zCF>eDRq%2TFx91InY=OEb}hR2x|7Vu{Qg-!!18O7aHe-isy7B3jYq8aKCm1VC+g`X zY)QrgG*Kg#>QJFw@$`We9(O~cAhJ;=}f*fXzEmfop%fo^L3oS_->BNQA4Gs{lKIn12yNrOCO|l6e$9 zQX-M?iP8MEy1LgP&XpM`S6zpxt6zTZoRdD%Zhm9G7ksimXX3!EsB6bOedc9Jqg+bD zg4GQh$@zN&$q_&ZHXaB(U$o0iVxYTBdf& z8+rkF;IxzV?c+LfegO;y%zK1x3T7OORF9CO(1N!QR4E|fl*r4@z}9E%UG_jQiBx~R zXt>y$?v+5fVpzSZpZ&cc-hLXv#xRt3s{6Yac$|@JtuZ+~ARb?@ zalsKDqrZZWBB*_R@3<7M)kCsTl|DS$cByW4tstO>kYtbJAa&G-wDah{MYGF2D4kU6 zur|jm^m=ePL1PbF1lz`Yydw_+BR&v<5G#z3I6hAU4@%ZL_N9n%nP}5$6Jf+#j9Bk-`WnVTd9$&my1QB)U z-fwpS##>BhbxKUvfu7S1WLUb`Re8SPSE>FQHq;YY1cR>`9wF3ef;WUwk+|pG%W!7> z*$2pirV~rK^v~PQew0$unsa>>Z&kRPK|fnu7=34WR9D`urlIt9eN?FN8meA%dk-l6}BT>>_T^P z&7kbnas>uOI<#G?DgHW0rtvTb2tAXTM_Tw^DQ%p~n%WCp%_ZkH78Y@!$II4lT8ES9 z0YcueL%`hXsr5A)8L$m8G0A=NDl<(J>ujcQia7o$!3uetm53cV zG(&1&@eadJd~<8cD7WkMTMcGHXdH&h+03mOa+pO3yEqv-=a1l(&BafifHZ1t^d1?^nD{bomcd9h?4c?L;N~pW!2E%~S}r z!(3n|HA+sbd#_w{@5<+NyBYkS|CuV-h|$L6%)6#8z38AW#h805U?8a<*>rkA$xN{Ihv=)8bYjl+=v$f)Ao{ zkwZZ1|4;?KqDH{Ml)cZ8Tj2wRruh1;_-_5f)DL<0dn)XR+=O$go&K3T7b zY;4kT_g3NEzg>{Ea6f{;0l$glp@8p%S&+we%fCnh;`sZzYkkVPZig0+!Y^-!dTNxz zEg@hFz91Y7hxxWsP|N+wAu%2m$5Q-4k3d)LDqO7;h-CfsrX@$&&-46P_bg$$$f2xq zY`FSB7?9W|E^_tk^=-;g{SPeJi?1NmbS0z+>rU-V7RA2^7wds{APyU;QT3#VF8gjv zJu$B050nYTN7uf9@f?XBHfazH_sU$|`9fl=mwUq0=?YosO)tZ4fAArQlKP7z@@|gj z61w(70@UM%cL<$fzkr8TnDI-Do72PA8@dtVsM+i4uS8Ae2j)Ed@)qQP83!7r{yOG5 z7yJ7>wsC}_sc}j*Z-%P~0}Tu~F!0ohWg5CqI^%9Uwj6~For5XVbD|<+0QHfj{4;g$ zpI0FY(@!|x>#GU5@#>0&BJk0D-wAL-*kNxAjZo_fY< zdIn@uxfay|&YCK#uYUsvKsQeOEC?lfY`JVUp#~ZA;o2ORtX1VH2H*ivnN^4Zo9kd= zAk-T^e!U$~Up`s*>p-;iU-OZ$QfjRzu%SuUWKq2^tRrHERaom9sc6g zZsXM++(j#9faFhe#5zOhzj0e(r6m;^l^|~7li(9=K-!r=&I~9$DFr1h9BKs;@xX&Z z_Xumsp9%l~`Rn4T$?m$U%0fmP@RZ#$v>)MF6k%#PMjn%MzeaidvMP z8v!;z?CKK88tKHa5v~@=5))M9NLpo~J@jP_aG)6KcuT(~GjeS zC>o9rM+~fb6?n5_3P^jzn}hS#tlA-IpC-KpGMtZKGZU(CBUVR9|9&<{aLQn< z#RFP2&LinB$biW76(R&ISlPEQ*Q*tCII4Z~8gxf%HtnBf;f<_cA@iZ*blO`;iVO9+ zC!>21WJR>1{U@es1JN2|$fw_&vlwV-N-Kv;@(uy9Vab5UD8(Dxr<7tUIt(y*^A!c5 z8Zaqbub4PgP?^#QOqpa)5#1W@u{P?rjsS_UaSN4s7l1quz}N84Kj?Dde+w8w-`9&i zaKlsF9R=$zzc?zG7<&MIAEJH?(?)VLh_`oiE46}Z$Q#7XCFOMgUE2WeaymmT!D9>m z9msbs-p*TdjHApNk!i4MTwfUP^klNnZ=f|w_*n8jb$_{@DnE~DmT{^|OeipxTx08& zds!L6^dMER)!-E|OnWH|z)I10YR z?)}4k@9k2lq27CAMX-tipN!jbf7&{i>u?T-C&7Ka^w^M#@smUbj}#KY0IiMrqBk1~YUBD3Gq|ZhxX= z2!Hl#e@OdJ0AdspqyQtqbQfNhs!{dKX_m%>yDna`1%r@H;)EL^c3VghN+nKpj=ux^x&6OTrSk*`F$FuGn~GzHq^vzKwWp zd@L0M+osasb`AR31AeC-#P$I#Uvd^d8gljAcZciOmyLNHOO4fc{ght;(%vd>u!k{; z=+YPDOJ_v4Dw|Ya41jf}-<-?8S>-}5koJ-GiI=opO(s2S(15j<$}?7X#d=EJd+ub> z+{^lit!-^%8K7M9K3psoaqGvy4Svr>+`4@Gs6^wA-$ z$K3tVtH!MBi*;)Z57EsUDWk0UlEz}z!m|vCx@ZIEkdW1}t|P41>*bdhlheiaYoVbD zy`}dfWdiL34>3BI-1MhodG|EV`W{QuL#!^*TgpCg^d>CpF-LM)2LtKvn8Uu> zmtMFO8weEao~evm(OAFqL0d5FCha<`?TaPpN;m(poYa<~iOeVFCe{>)v@%Eb?UdN8 zm@@iTXb>iFu2Ujz;7nzT$f~R+?N>Bq%}0eM?eqVU5!JV7hnO(?q`FAveMiB>$;sQ9Zs3a3T2TJLP0go*F zDrFAWJGJwF%YOcw`i3McLVdFhIY1BwbLSxU%d!8+oy8N1IN3ot!2N^?erc)?3QQd6 zP{M)u$(PN!gh|R`CgH!+fJYp1w5{az;_^BPAm4bqo&-6_gE8qA7&4E(bV7vu=AFVc zP2iYY9<{!tS$aw0b8qG>|5rAUZ7B1ZJifXCDK0z(zp-8vPU;SqUm6Hy-HD|oA8{>0 zL>U;c`=cQn#35pKJJzrm|3l&V%Ht#dnHSs`c^iOjC!q5%t_@La1=^7!2zbf>o7kuk0#vG56?AhEA~7>2ljwdlu%G zSRw23-0jzzw6cwn5v6vau!!7mfL|Z6+W@8|@B#9);%vNtQdys`buFMEzsVA9uKZ|b z@;`}m7kF~%Cu04;@` zfW>sTkCQvaHVdYVZd;muv617@btoX=tD}#=y`U5{{T6aH1AFCU`i};SSAmS8sWU0p zq#qPA!e(+YueBi zi48boSfB?tb!fN6=du!-_VCo>#Ra10t%MizVX0XC3u!<$d?kKpPU&#rt zU!;9Bfld}}C=CXc*w_9?LG0`|fX!WL>Y&aD??D`YhtNecEieQe{;mesL*NDqi`GB8 zn@U8NXLD>z0RlV@J^?=%D1TL|-TK!)f^eh~Zq?dN5|g$QbhJ7NKI#A_?t?#`g}D|uly)FbK%BBO*?{q!%UET z_rm*8`90Ww30U+sW{MhQLC6fV>kP)qSn+zQ>ma1vl=A?0xw`)t2<~>%@J1-eZx9P) zzKdD6g$*gAzq0@V@y?w)SaMa~uYqyl&p!V=pZ9mR%Wq*a;pRw8BI!!w@z9f;hJ!y) z{1oefbp-XoPu|0+MH-!#r_p{_lNb*6Q#>Dpwjl`PJsdJ!5DC|9)FN4}kV5>%Q$K~YF@uYn< zT`msY)UcO8Gc|PU<}LAxzv)R}^3SuDZ$LFx!3|Kdg7lt4tdPozagCU^S>ub36oYWc z*w={xu@7R(C6ur}SmoI~pDg~oI$G-IL)WRGmGt4`Cw}h)W|rk%kbv2-r%*sQZQQ{q zGe_~H4$p~9iszJUDC=pEKFI6fctRBCJn+HpPA&u;8%d8Qu}U|(qN2N5 zm0JNIYt1B^h@o*PEZ*I`b%M}CoxIt!s~~u0FDby?4RXc>7IvG$Jh>E_J7k;Viu?cGAKyGai1ps7;sX+{KLMxd=~Fygyf!Nke)MWT?tuDfbH&(O4ni&3rE8!GmIVmg29yLYYjR_ z4}5+5c>4RAk;yz^;Qd4dlKX%!IzU#3<@u5_8`CRc&me7`;UwZ9K+#~q~&Pv$# z%sl``B1%ZQ2xDSUt9vX+A~N4>edh__OP^GPo953nyBm<6CR>Q>?o;zUkkPxE)ITj! zJWN@ep+skl_i1`+vO!1>cFVq*-yz`1r+A|=@jhHNt-B$h+{us{uB{~%5qBZhwATg(|6i?+=P-|H81JAF9F=x#vU6`_O zN!?IC54~-eB%eO^Mahd(HhHz~0}q2*>d}-J%X!=@K3kiLs3npjaSSv76{RXXS1FNN89tH49)}r zSy%s%Ak_!uunIj!Jx}<{LVh0o?ZhL7HXmzQtwVI;%>*w%9WUhGGfpE-XVmEhxm_F7 z3gzra4SfZh+d9xTS}YxxCxIf`$BSlRcIOb##&GA>$ZOLk5>+|kz5{Ei`DzrO^ zuCt_3Yt(!Z6U+Nto;>_nZ+6|eJF39cePpm4YtRzo|CV}PmiQxlY;ci%;QKMo;}o+EY*t0*$pmKK8SU5d+il* zJdx^Y#Ipm})Pmj*L4@6b14NqcKEb98=i~Iwr8VVknoCPLEc=tpB|nx(V2|=C<6^D% z$EM13)DsWQ^hTjjd@$aHZQ0Xu2uAywkyz8)_J#^`6CoCfum2!?XPQYgC#E(ZKa+ZS zbAxg9)y4_Ayv~JOtKx;EejpQTN8_6#ULB`Vb~heD$&CbVbgmdYbY$X&JyPx?$*kIc00e5--+A<(xYZyoH1sm6;UtwE# zEW+$7zQ+tZShU+5D?S8uGc9GD2Fbt!ewE+OwR+=j^5@zSl0#65kIgkm&w?jDYQ2dh z1c22?!JpAXS!fipQ%}}|oqAriG{y^zgj+J9WpgiC(ek;ME}D@yp1s0Y@S2IdkHn*v zR!ZA$%}|s3;W^=hOT{+k#gzP5&pL=yYZzE&?qJA+IxLE$gc!()!SLx6kzkpl^&b^E zC=LOXFZH1Tjk7y8QjTwxtHm``ki=vbk_Qx_a@!;642Y5Ji!uE5(NGpFqc6htEH~)z>Skt&(1v_ zlFbl`&dzPwe0my`fNHFUObb{5^s6_JH`hNo*TPAq?I(79&%o)XmDgnzaxD$$wQ}$) zv~#M(8LKzU%oSv-5FaoKNNH%iIe}3+>6Dv$Y)0rlLz)>f6oMiC4ENGRMGC^r$$uUD zZ2_>2aqL!x-1>@?#jx+i*RCB5E}42_YA|5O*&o08+QH}QcSiWM69MrbwK-WdTsdlA zm{KYbCQPe*0TGEeem8C6y(SeSOpR6X9H0iPfZEx>(;+e8A(E%)Le;!JPl))}&xI?i z(sfYQF>!s$h=<&>97PjvjV}c@ zP7CRuG6W8CPh$af#?3R*)i)b2IaP)hrLOT8^HSIL-VsAxMs!b+4}n4Aq0;K8z0A*j}=~}w(m0fr(ta6P+{f=V4?UHzZh%9tGRYx|%AUg5~7e}E`mWtq}?hfoJ zm^h{PbNbiMOX3^vAG}h1wZXx=n%H1)Ew9s>em$v7rC;t^TOsZ=)qje@d1}VyaPwv2 zVOJuKUNq<--yG~&cVR1U7@VVs;%6u4pTKzRY7xuf$VnD0y_mQ0Te?%lL7(73M>7&` z%O5&q&XaUX0OiHBK?M3kDK^m7%@Iikwi?9Zj)wp+J8w5#J3lEHSa371Q-_ZBWKq_5 zmMmynxt20=4hocv(~Pj~ICbTOuH+naxwe4hRO4s9l@yL|NS(m)qWgAm*t~f%DUIV> zIxL{xkhknM?(G~oQLJObCM3~y2>U175C*1f7fl3Qb27h-}C(-os6WOaFHv{q zrs|-FIMX+)cVFlks=Rwk(eMwR$7q^`^V9I8IuySM;wVqL+rHdKV9Wa~eEa~%J>JwO zYm5)9mUCjn5BIr$#sD$hV?QnpQ=5-_Z;C@O641z-=NKhUkiQmx>n8PlqCO$`#zebj z1zn+Ot=8sKmB*k)xKqxcMub!DphhHvi~j-IqjA92WH_QLd2srJyhp2P{BXn?^r+AM z>XIk7I=p~r+F>JZG!an7aYE;=BwLO<^l049tf%6-r&VT>RS&4S4Q^dzEydifi7Sgy zNfMMK2Hk(Lb}XllN8dQ!_|2K{6Akwzha_JdVf2Bn35)TZxxW+?DIqF&P2I_-BabM- zJPKBnwGzNe1T1BUZM*2D)pdw@_0j6?&_?au?-+ORlHoHSbB1V&Tizj}5lj5Qv1{cP z*L1gw9%*fjhS?`{4Qw*J?AP@C9RxXi6p+!QF%t&s8d{Id|<96=G=}{mVZe z*AzWCBlX<(LHVuFm741-21qkb<3ZqW)v2b&au5Ex{ff%DjVJ~DxzZg_GsK-%%oM~4 zWG$!2;y;E5pP1mrb!ByEocbWy{8PI^PA>^;+JGIIaQZdMQrunIc^4rkWj!0J9~`{o zGM;@!{>M-qU*e^Mp0hqj_eiuHqMn9VrO*YGnysFfiV^AHfu62h0x1% z*|x^>=@7rTDXnr3qX_LyVLs#UICAPFi-*0O>h>?~E_u5-Hjcb0zeD|eoQZ=WQd5mt zsh3adhG`;zga{aT&pkd@3ZJ;(1pctx@p4bsaWhYO#oTG5G?jhvBwIkm4~gOx2ys<0 zmKsk~zFXT+MLfk={u0mmpQpc&eSYhKCEih;8tIEIwjEc=alz>}s~Fw#GD~v0tcLG05od2lo8MEsMp?VUtDA`k{4=#7 zEGq>nu~>F3!0(7q;N`fC6)uVJ9=l!r7F!hjdN zvZV;!*)y%I@2PHA&C8>3C+R2`NgPSjZlO!GD#VVUI~1|^d2$59N-4UnD!RSXS6WEA z;zW&b?o8Wrsb`}pgfBYej?)w`!J#z%Wwh|q!Al|8q$?UH)j+KRblW7VhqxrGTn)E! z&D;+A-?lM+LU-d4x~teKXHj7G{g05|Gc(;HOIGBdbb;D1nft-E6LW(##uMxjTse+R zVA%=OqCl!E!BkC77ujiBKzY_uT#~4-w|(OgL(IUX8yg&s?qLrfIcx9fbxrt(JI1WyOF)HMv*v_7WMz`&mpXu{*_Q0Q{Ln(8-kOArR(rHuZ8e4|{ctEFjYQ2}3&Ll%mKp;YobW z67)kU(xhZ6aJ{7OoY!h;M3NXX6AHKJRpO*QIxMn0?F5sBz56nXSHV4=S0WqPis_7T(p zJw9sFSDfCmj)ALPs+Q92SN>zkn+qn78JXNYEL?|EUD2UMVjmz=jXOMzK%;^<{znAY zj`lpu+2^DLDW=`vwD`i3s7Ow=vE=hNW!0PNm!+au;hoi4m}&o7tMH*8IN}TR$A&f4 zXLlBsYjCF@pF_N%l3$FA5<47ovnr!Q zv-^eTM{mmhU8qOiKuh~0O)-k1JFLS=Yaoo=cQR^FqwYVR&3`tU zC;GcetNiZbX{6*aW86;RGR7wfXwIZc5!29%hh~HfkNJ%h*_oiH9^oX*aHV-JX)asC ze6>7P?PS9@Trmmv>+Q#QhEmteF;EB>NSA^%Eu>E$MZH91qoiKX3C>r71TP3 z-~qUFCxnwsFhMV}Q4^-9MZ5FHQi9^gxpsT9v~_o*i2B$RYg#r}!7=TdAISm5=k%X> zt|~Nl&Y){LGYr2KaEUN36O+*A*16bJ?wMh|)By zW0NBwL29}E0~?mA`cP^p26t`0sg1j!WNtw;#z~Td7HW)*%NxqdwQ9* z*N)@VkImR}E_NR5T< zRxa)MLI0!Bz*DQw&G2vcSuFunAhqqNQL*7GP*?#s84foN)L#!f7(#W^$35>+-gcg^ zz@kZr5VB5qZcsIvH8);u_Rm&7d)DTLAfd6y--fYg^PtL;2eY1Thw%9V08Ra+c7Hd> z;TjZogy(b)Vh?jP+PKvE15A)*-|I|0KSDf$a!B6Oc#Z05YvS)`(2w~jFd8nfZ*#&G zB!wQd^pMo{`-b4tH?6HK$&?$up&k|49s3RAC`veMvLPz7R;9^0wK^0q9or;+MhM6q z#7gOni;{s)IB`2G=;ud~Acd-_haO_lNy&Jg;sxrLhC+7!nJQVorYzFOJkt9-becf# zcm`B#PXnxScu+%%qX-rDA}h!m(`qT?L3;4qDvQFkWKGt(LOZ)OZQq@~+Kn)iQIqVg z2y)Ymc$)6WOX|opXs6^WV%4^z3RlCq=~TtJ7!`qi5;=#+f9TOgi}D>ltv#J^RK98; zta+pf3R8rc@1N2YGHtPUg4a%D0`6tv5m4-ZO1rq{4mk^BuAXr$k@48l9BPmQmaMQp zz~G7KP3_{uw-HVc90h8CcCOYzM0mQY6R6mftw}lI7rnHPs@S+`N&oy^L`Ic(C0S~!T()O^{uG~Zd(Ms>lt}qgw$e4gX14FmN2f;6PGzmSy1W@Nq-AaHfJDl1 zXR2;Rx%^R|56@kcG-ggtj~@k1O+5O2TH)Wv4v}B{iF^HFj^OhYsIO;Ahds|1ud9ET z+j^>lUM@bG=ZZGVCM{~1V(Pf_)@kUDI)-d9Yv!YKz<@GdosdE0!mV*zWF?9ekgOzx zAEU8sL$dul8x7>Y2QLhQ;Bv3a>Y?M{!ZEwCh@|=O4U7&EM+>6S6*W)wP3w0!_Pf>@ z=WWO>v49Ii9rxQ~@%A*1o~ZIwqykhWw%vKq(d1TvVJx?1v}yRh?rR##7blFQ;gTMA zEb+&q3s)2E86L42uBlGKsH9jI(jENjxr*N%{;Jize59GW98U?7BKebJ2ekjnXYa2R%M6K`^QU1%e zciP|%)E6IvZftMR1=OVsj>0mD-;SeaW?6E#bnsAk?EA5g9TiCe5xF(PzL z43_6u@+uGy>+Kse&>!^UfTOuB59ka0-1IFuhxV_zq}HC|w%$1ru)N8Zm6D4CKsxj7 zO;ed{ylXhlkoMDr#}%nMagPthllF`{hB5}kqp8W_rQF4zvTNIJVHh8D63>2m8=*15_z8Fc1MW1(m3pW9FT7rV|4Al{w$n2}-m zXK-<25$lqzT5*}V<56)3A%u?MmyAjnB;}zW7R-CjYFpL)95C6i0WKB)LnV%&p;jD5 zY{R=v?$3iY++a)VG-W^PxCNCi6T8^Wqr6#Z&f9B2}J_v!rKM%mZN4 z+3SQxZ_&4TFSO)j^Hpfcqjvz}<5d$>gJCKkCebq7z$?k=hJ4(zWif@;3tZF~PdLLH zsDQUk1u8Udc(kk0*)x?_-?{(MvR=iV`6{ixU9dXeC>eVf{$x1f5N zyTtH-wFri2_!&=4ZGkP9?d#tL~6v8mMO4`Exq@cFj8O zYr8o}ceoMckcX4(=pfr$6CAH(g7YB4c3VDY8}NV4A@{)i3O%QMLo3Y|*uCU)TJY#& zMloxURdM77rW3<@XFVuM4A9%Ih&lR$TjZ%|vmCp~>n9QJ8=N|QtYjVrHz*IR9Rm}jdqFE$sUK&eYN>0iWxpixhTe9S7YVH8#k#3I zxE(@kkL?a{(fS)0%ZC5aUln*G_Vv1F6CI&c)9z6KSC}48VQoJn|RA1x9fClZGbYhHZ zgKKhh-+X8*YyyhAy;Hvx*wnBlHb38skdr#=5PWomOMkZ$P$G>1d7$-AKYj&?J)w_1 zzZ%o01xTD&N8ETh*^W7$6%8AeXQkN|qyH6K}U3 zoEwZa5G5{bybJv^jp({XeD`jv3p^lF{|*d;+t@2c(8c#7556__ zJ(|&1z4l6D9t=9LE-a9qj&|Il?eGw)+@Y2wqFF9t1f*WAup~BOjk*ikQ$QR1)|QTx z&0q=W6kVW$&20mE(Ey8AkE7Vm!~t7qod@N{=Fo#%*P1asi~BCf6P7QP7 zX+QqB|6lj-ziDnv2JI^C(j1*oL9NbfGZnPj5IV8vlz1x$5wd794;)beH+{U7R)a)d z9aU%CLpgTEI$DuzL!~TNoT85%p1ubPtpm3P-n*Uc)4Rc&VtI?~Z%4u4M#cVB@>0^%)Ii#H}x*ihf@@~c}DZg%j~BCM`0gZB93PPn zM&1m!VGgz=qNj71!Y#!v@!T!VIn4f^{sN5lD&`p`*g|ijAkoWHzqho!U{C`y@Ob2w zeUD~Yq~D-V!K$A{h@E=qw>NjjXrjEfZ2aIb)p~jByr-Xm;81Gvd9d(pkeYg*i8sJq z!}z$Br&cVmN**gO^S0>5ZypV&?3o#PS#_`(bMR4{BGKyf0sKtmGE<1TfrMr~xS)<7 zX)LT(>;*Ce%b!6yYc%;9$!R_sjvY3pqli{QDzfTtuKLwmXlW%oL-u*;1Nbi8cyk&` zN76G8KX5;$Z6Bv(?pzfgJy+=U`q5 zozd)_s~%yfzO%V4Vm3rFT@yCSg2%|axhC)SXzA=b)_n6FcDLOzF3lFD$1-}SrH^

CrBnQ(wfB*)AKwui#wX!0EjA6TEZ~wSI~s5q;+^^V~mrDR)4X(XbTOo8{C_%J+?@vGBx$~ zjC2xIWYYs|W4;^2|FyF3P<)0$yWHz$xiPf}mFX?|aLa9gi!0mDY0VHqg_T;K9&95p zveZ?qvMYN}f}25|md3LV+j$JW>X`juT5Y$dz?STGcD%$WSK*^0$x#m8*y?qz4JLzn~QXaY95?;vuXG zIIGz-WM5=w<6w!`5VtJG-_Wpz**uTW!IKHw{z|+MFUd8e6E*3e0$32oS7EY1mKAA- z@vFY}?yHCt%kBbmrBu2L>kV2cl&TcRE7RLM+{HLL!JL`< zuTGi}@6V@D@v->QIA5;l8|sUp*V{sMLPGKgWBl(Tp-2qKOI)A|bB8MC2zLJla6lyO zq$!UZ&Nnv8fm?E`QcNjvD&~8{2r2U@WW9mfrhV}6PuL49D?WF=44p;q`A0>jjNy&Pb7QSq@E}H|4nsIDC}|Lt3Q; zxZjG*S&Wh=XM$>J>?naD1m^VsM+)wD7)*J!J zoyoAs5B?plin{oZfWZ)U z5r*TWtI43dbOBMl0jb(qyVIFKtsf_a--Lposh^s7C@F2{P;7NdC#G1V3Lv6cPH@~a=`jrsGH(2iAG)yT9Afm8q{xxzb`1_ zxY^$i8R$w*Lgo({5w`_@7WX{|-HUn+3m(WTY&Yyy?|K$J2~NIP8Tp<2ArNH!EhwpB zkgvUg7J*0{x5;DUXh2YS33dl~uwj6%`tLVJ-hAC?yx6QTf|~sO1*lAg-*<`n#<^}o z@cpWYE))4=*jp%RTYvpOd;8soLdOE;Pl8FsbJ4JU48FjmErGo~V3Le45bN8&$uoOA z+qV8UfB#(`uhZ{e7p9wqIVEj!AI|I3MhoUCa+ten|&{g?m==amD z&jx3o7NA?Sa00p95=O%!>y63q%3*{Y%qjadJ-L6p*_8P3pH~MOcGB;sCzluRbQvD! zz5XJBvvm}aX`#q}8vtOq_uo$ys_7pXp0XO9=qm042U);|iGRGZPvEM9KY+`BR*ZMj zSpIT*Tbs|n=TIa6mY2xyk>O3928#xcdHlPr;MX60z@?@sf@3}4EWuU!+^_%t_n#m3 zWBWC}gWGij)4cz_-{Z>vKYz{NZUbB{PS5JDHs}|_pc^_2d9M8Z1*itlVG)5SnvwsL zO921bU2)yHK(1X84$>+6^+q|3-$wC2+T?#d>i-e@?|*ORmoxugn?e1rMgFf5{NFT! zD@ow2M$oKFn(VDmNJ&ciJ_0)p?Z8ReFuKq!exVN-@HxPNDy^_%*cOc8Ys$AO&b7$? z_4oXK&GMMT_$b{ymrT2A0{!4Y1IR0ECL07>FbZw5b+1|cMVxPv<=yX}$&OtFZi#>q zlFyG%(r6k5qxdUC41o|+3&-Pb6Cv>e!to36B^MR%q!TVi44>7f9lj3QyjB%lMn*`W^OG2gOl}45ai~VO z{CUpqmJ1#)&bl0G@~3322GQvkzTrOu`;6V>wdrD@W1-W)W~U zS;W6j)PVn{Q}er7NA zoW?GLuBFs%CUiMBml3*L8j^}4|I78~&@FNaIPoN}nGpd0+BDk-c;w2$TnrT7g5E3LbOAuawEFvsVKa{~ud}A=Sjq7g7(mnl1@8)B3{X(KNT74{eSGVzu)*jE*1En3q|;E zLO}fO{i7NoHD3(t+fSSzJP4q_CNQu^9UM~16aZN~n4`J4=vn1AJv9P$H1t5TZ9B$* ze*sD*x2K=|WPDOkL*h|e99fIAEwaw=s4A((G|8D`;VL^eHP{MSmy+8KFOmyt7B7@K zx25;9%E%cHGGt4wgdT0kj~`rhnh8;73bLZ2m3No9B+retjg>u@0Fw#$re_>Q_Ob@% z$2Rht)0c~`4)zYqbqiUC_cZzrmMbh~%&_{~+yB4z-ZQG{t?L&>-RjwDm5r2(mMe{QCg%%S^|W?Sz+(@dG9^<l7S|gPf8C_?5Agn(sjDy;YVl zW~ThysNxDgJ_{j5gMNC`0cGKicd%7ET(LPWxOD1C=FL;rtK7!QKvGX2%G&N?v&41r zGhSEQ0q{I!EpBFZeV2{1TjF)pj#rPG%uat9k$cAXScCHPTjjX=GZE+G#jnb z_=5zRM)a2H``T`4ERj$F;TIs)NeRd`IqXjDZA$;BBho`b5D-d5@|&*+OqOLVw}EE3t&iWl3Q%zO zqbLbEJo`QA>SMNbsy?==L$c`$lxjpD`-p6;O*UOe)P>+jGuqalxYdnfAD!l&+3`ku zYsRy5T61Ib(djz@^i0?l%i*tyvg#*I&ZDP=h@KIlI)J}O{_8}rG3zg*l7D+$daz!( zKBY`@_La#*aUt`I>$a)X+~z_ZDDm7?@xN^Q zs_KamFQ&sz(EE~o))&h6eb?K>Di1tX>n2)&mo|qwz_z-P0Jxj*-9*3x*b=@RK z4DO}cX%53&_Y!RMSwr#u1AO2BA(4OGZE4k*&Cj8Oqe#uDvKG_iwbjxx8d?g?ZaQEZkG$jpANB8xzLwu5{9bR zy=>S~$Wlq^7B!f{wFY*}#P@P+33o8r9-1EnN%7Kkej3GANS_)cjq0xmbiPRqvtPTz znaENF76gMAHI2Kby&*_EzYEP-M(|dw+sk$$e@ReDvV9C!o>2KGvl{A1g31c zll7+@(&#vIg67w31*YiQZu@@D>XEPx{JW+-)3wr&jWR5Y5IHD@oQQnoZOX?03e`t* z_LLjelnd*pVPoe`Ly1NQ$|xSKz-g{GDYs`FMH-bOBIw+&Cmvt6xWRyZJ--=Gb+8v9 z*6oFAGY)6{%sl(o-N0~jg8vDptdNkgN1tj|(MXr8-VfIm;$h=nZs#mPcWhtm%yWJn z+twFqyK|%Br*913eC386Q%V_ZM8*H$rX+fKz~mur>Zv>iMZ2~xgl1W zJ&koC6Zlfz^WsGt>rXg|!+N3&q*P>PeW~lLELt)9&g6^@y(lZT;nH)N1ThoQscT|J zcstGJ_*0#QoI_M)@pHf^Gh;9|+2XL_T6mrtvhAr=eY)wJsf|8mFu;eMGQooq`d!#@ z;!+I%tM7E#j~WW9dx|WYuM@dvb1Wq`;brvV9y1LAE`CjqJ0x3~>(!PlUh7kpQDf(3 zgIP05rOLIsoD4!5e2(A3+x}fS+B+-~uX+_dDz=?^tWI%sTwM8aIHiW9r2g&#$z8>CS%unK!l1=KPCxkmARiIs2HbxflKKh)vyS<7gdX zenvN#tN6yLtDtX0-)H{jtCBRjHp{@b#jOtf?_c1NUGYX&e&y?QmqJcxMd&5ica%4^ zc)4K4&Q-gOdH-Cttx5`egWZsOEfXDl>@9O^qAiALvHI}3^$!M0$^H7oXuHl)qS{aX zj^|z>a=%O}YoB;yHZ#`Cg$C%d^_5jjInCR`vnMzy)v|uD!#3{+UN5ZDO-os&3IHq) zl4NrkgsviQ4W38|*fb6a%?W5`J^ct>Jd;_tFTL&{b>22j^+c{lnj(zx$-JOO*w>kWS0>8GaANBV=xM`z(57A)4r}=fo%F3yWImbE zmfJN~)Mg^#^!2R!Y_HDFvzM`G(yP;t_qLKcuxTjaumt<(8WLdM<}NiX-h3ullwa9e zv}H>(9)wKkB44rH5V3wzHriMbl4hb^FvJD5>ot6ws~CLO?PiqS?D524VA1cpLN_Nei;@L%T@9Lgu`lLcAfjtNkpXMy~;`6J~@UrVOy`gY6)Q*I2nO0TX- zm@s4%!kuRRcJDLHrJF6P*A-^)oUjcaANi|v$%wIJ{3QSF-4=@#*3J_QUi56(_g*1v zCsQ_xEY!eQ86j@}HszbheI@V>ysE%z0n?UOc%g-`<@d9F8=0!o=4g6|o%798X%R=^2u6*$r_?qZGm2XJldo`&z;arxn6Hq`K(b3*m_DG*(K@ zhL<@Lo!zwIL9Jjvy^oGhO{e7W*}o{=q?5n_`}g~NHCo@EZ+T&YVbIzQR-HM=+mv%% zCkUw%!2Ig8=B3u84L{aYQEriiqmijQ(m z6lj$=Te(g!y5i1m$InsRc?K0--_JLrG+?t+E-%`;&`l2_?7n)V_(&SBPU&?L=k1(( zTMZrW6X&(|`gV3Ng-1#i^jWk{!mC|_IeDpmP&YzJ#SeYvh#Ntsrq zAmia%Q4mjNdjU)`i645(4^hEB6C4O4!d?+N4=^S57Z7}fd0 zFM;$lzZ%{?d5QjwFHExed~Ez#8HxO3pC@cFas^ph`7E`aGXa&;1Qtf(W@ojO0_*44 z0&KQ+ThN_U`{YrTM~Qxq;y*0|<)m(tbRuOV!9`o=SAzLhuQ%dN6+&XS^&KDHG|cmS z1V~~8={8;IEHIuH!nfGjW~|H65Jo z2R5zOm5Ta$E7jzSv`yK9O@Xi7&irKF;WH49J+`&;qq=IgzpG1fDMM`EfhjRlU||z$ zq$%cm_WH2BGK}%r@UaNwDVwwkBe;+ll*!qwUNo^GwO%(mG{mb;aEXTX)fz800V9uV zcL+5X^gs+ItIr2ZJg%7l7EDpj(z;WZ;ixJHAn8VsmxbJlmxLDRe5EPJ=edec`{(}7 zucwbY_ShyZ0mtuIhPgiN(bpw8Kl7M3LW(Y~BqX+Qt^IH*p+>D8cmYWcAQ;m@BwX_f z8GH8j+Xm#F#B+1wK<>#oKRWFeFWZ`U25gF6g-4dzFHy1Gnxu_Z5tJ$8JJTppQh%#E z&#Bihs^T3TP<1y#wi^Nl+3VnDa40pfQgE>1^9g|D1*8Sd^Rc3G7?+)6s20*aYe|ry z6Gn!-@mS$lJ})0Xouux?&sSw+`3g6b)l?zKZQho7C;e-rcfWWPD-3yehSyZ0(|9sI z`X=$qY}Cn&6Kn#SBAxN%^q4R#;#EsyRegGG=S^uL=n493B;9$tI$1*8{_k`qH`X+(2Eyhb(#5) zFWe610AT~>3&XQHza+VQk|)Jz>M4&o9G;0RN5 zxc|pE3#DZvjGv^f?bvf&xc4|Cw%}Ed9ULYeGqlmQz8N2_l-(uu%Mrr*G;8!+AKZoA zvkz?#$UHgb)bWUke35S%9TFD3%N5@)D+A-@se{{&RNln_jT{oHOm^$%%oJE4{RqZP_xl=aa3jxu*dce322Gskg5QwbiQG z=R(56_>z`>JX^^&ptOwUiDC~pGK)eyB%}KGy4Sc{lml!R z?!qw3T}6kSTCpOj=N5?2x-Z%xMWYT(44o%NV`R?UqB9yEJLK_-WM`#gq2!DbjWD@A z>lWLx(gV-jgh+YAL?x}2yNe*B`y7ru4_SI5F}(eXRkQx*vsQL-jTvo=dbH;1pgxq* z2DO-cw=zJ=jgDU}?)D@Sa#bH3X=L_eP0?YhN!;u}Yiw=M&8fnhV`1^mD~>O!T5ox{ zWDMS*l61YgOv#E4^_WtH<;QUu3K2J`4h6?d3fM|fEL(egq<0chZd@olPQ%pqFUsR= z$X#QgZz~|aHhKcDTA)>5^uGl8K8o0=*q)zIoI6IZ< zQOPc9i-fzcX`S~>k}TZ)xQPe)FoJ4=@k!Q;g;H(e$TO(bAq{~0e1ld5^}_bt57}Qr z42P$Z_`Ph4>oZQ+d=_MclhZGqD5je&4w_h|Dv7(>W`3#9a4l@HoN{=Aj^%0cj-cxL zX3S+dB+-K_`|s=75y<{Of21FQ6dnD3%MtAS-&Y^||GfSGCS}g0$@%ar036S(K=P#j zX6cn4v>jH|ZMg-?xO*?}-h|3L#g#8z-nW1(MH~M!wl-N89JpGEdwWnri{;N@xxqow zM%mHM^BgC@v!wyhy3T`9Mx7x1S7uq~xBWoIw{Mi z4??~JfD#-7F}~n8De_rQ4JJq+j@Y3>gSPw!M1t=5DKIjk!;}^)=he>xJX#)-wIi9& zYoM*@r&q$Bv;Rd6OrG65D*u=TRzCP-=8MOjl4`WjHVtCaT+;w#)tv{{Ef0o@ zHJTh;2OzJFj{D*1$=3gbIzNQd!vAqcP!+UocSwGc*KD_yiUuR|#2j7=h4&uKax%4m zOd5)gsm|DecQ?)byH%UfGA&=fWI*Hjj2bW}nL+|JA3uUg=3zqiAc%R^BY0&Ly-!OR zHSzxUN4fw?DC>minQK%C65<20TT@kqC!+xjp)kJI*Of>BQ?<+uCrKdl5u{=_aRJs1 zHyEjaFp>j9qIrBB-w44nO81rQa+7OkTxUsYfD$p`4xa?);=_C5cg8xc4Q_(C_b$sG zEJOIp?9Sc{xgLxr1mpHvd$~0AXgQjHc;y4VIUFT=Vu;iw2|$V)vQ&42w=~_#+x72g zS8=JO2(CbKf&JQ{G;jYgD^1FYB6L)Mqr@sKe6FyEXe3&815oK_iyix$q@mRgfSqYs=6GS=nOV>DmV9WbFX-nIiDmDOc?G6~Ba9`Y53 zjl4-!^~&C_eB}^{VsV%u3Hd(Y8;$_F?^wLS_M$zADwk+|2N`{GFQ3m=>`rs^=rJQo( z1qb3a8POGb7;kM~{40lRC>H;{gqx=l9@r*lZV>PgzmaYkwtGH8Vy?Pe&{>0z>$C>_;Po4 z_n-v5UDsvNYet9HM+j8U(SS&fKep2wukdY{~1#;bJ z*c~TgHxzR{K{XJ$AB-&$@{Lf_WRAYu^}HOUtxrl-=14k`;J~WQYBD@HSWkkGI8mZT z%w%h>dHa0*Xa%qU4IcPto#C=B+wySLv8!je^&aN4NU)?1XM33Dn)Nv8J3&`js$oQ9 zqe;J1srsA;d|u)4DzpR?Sy|ZyO>Y>VyeBeP;fgES7Z$6EFVhNaOZ=XREwc!MW=JC( zef1kHX%l8^j|RB_8O9IfUXomn#Y<>QB)^a^2q$nn2n=3VS}K4`6%^pD+~9G!1%PK? zTsvqOm?3`{GuJQgR{g|p52P6=l&{tlff*T5t(qNb;58r}{@m-IQn$XQ&-Td|%iqLz zvZ9&v3tsD2LktjH`~xkSwpjftrTg2#o>TRKUoo?rvoio+y&->Y$n#Mq5UA)7{XmVP zokMxe>qd?<>_>=AbkB~ad;t1u3U$7Ak(i_DJI8|}KoNNbt++Ak0()2i(k=L-3yRmF zMTk|M2Z*0;OwhbqJ8HjN?Wd5D6MHU%vHW|y=g~~g3{#8Ubyt9$DebOY0*E0l^sf{$ zaWT!UoqOmalnXe>MedrzZPaoUz-fP4CunYHq{*obc3xn}(E4GcNMZd@PoJIDi#s8F zVTjrfE;hzzcR4%L(0DsosDXBfq5rB`NF{{PTIClB&=$x&v(Lq2V`KYLuM%KF>(qh* z6-Oq8@qto9hVu4!l`G}lEk#nA%$ny~i3T}qDGCW|E zYR?Z=Nb>u4g$ku|Ew|DF57UMDL~E}tGg{cDgHJxb|fTND)3Juyj@|-tcHePpai$BcHm)H_dYizn1+%v z^$f`T0rVjrO(AZunI_(U4;WZa?EHGXP;U(tVk2cujqd{(AI%vU6-(!A2_6 zH>fuPSFTA6_Uk)T9Fd#WO%7TKnm`6u&j*_@rwtPlf;F!P(( zs&G{n)INDP{>-+!sS^u{^(8EL08gag=uIIvzbwBJStCeiP zPqnYG7iYfxIg|c(fQ0PJXX=m-#WEkCUm8CVH`xD5ibbh8=c?1aYj&LK3g@`41&QZG z#a;4l>5Jk@HEZY1bvJx?N#Uvw8c%LNa8Nci4K5~*JF4`CAsVd2b|M8RrYfYpl8jAi z2+9Aeiy$1Z#L>$}n(042_vgzI|0!rSIoJ9hcfIJX?NJ8{?E{bQqG@)+(W~&23GHgp zUeAXg9^C@k?@j;nctC~!Ar2tI|2gL*e$x%#8(hU$Ehb|2=SkKy2BQ<#r7>YqA1h|3 zHE|Zyg-(Z!6Yr_)^>=QnY<$0l*X`}WsvN9$*G+AgbSDTKS7o@#^eT}j7J!8{T2D|S zjwzBx?eMx{iM8CM4=b=GR32hnasPL7>0Vl5{mu*w)HpHfDx#@dH|8fi@F}ZrI|~ri zP+}Ol4v*>Tlt%ZY91wkakYvy-iW14CgD9!d3NLv!|Y{+WXG7eMqd{9%6iw z);QrhzMgix!-_VapmTGwr=E3P*?`yz-F$>=xG_*=a3$pTSGCP9ZEO5e{aQ5EgrN4Q z_h)Ff@}|?Y@^%ejDC6W(mDR|dAUe99MB1Gn-79oF#4m^;_C|`o?7kwpsVHd|U#n>y zE;OzVUo>E3aQlS!KWOx+d4kGx~{Lyd>4s@k0Mi4}DRZJM#ab?|9e zuV6PhYhjH2fL_``KOJ+ij?9kx_{y|_oa{=yJ5 z$)Im$Loxn6v)hkZZ;Ob_ZpSov7hYw%3LFM6E0Zrn#UW#aAyr5mPy z^ea#$)}r&C>6VaVS1_Sk4R28*3KG5$q;{%-%XGPZM({wG>%)7=8xtq!nmB24yGUO> zDsZ#GlSAYqI}A5XZR=(|*wDn+h+#&c7S}Mct6bj>>S6XH`D>}jyY0cuU%&7U4j^2E zEbnr?SOgV)c-@*6N$;*y<&bHo7hNXxYgBH$;ub;IgV31vwpCc^q-SMQ?gyFAz4cI7 zK=+sXt&;B^wBY*h1R=!WI^mJjfHUOOyG)&0&Ye?^yNEzP*Ip$k6TJ9##8tcuhb0&o zAxG0RCh60MOVm`{)b88CnA$2*Og}yujl8UACM&7%725%Is{6F^_U{Ep%szfKRx~3B zLKs%?dHdG^iHLB$mF~Y-c&4{gh^2Fg8Xqd|xVW`K;9>0)FHt6247rY)qR0^f8TAuy zvgqj2M2_gFZdG?!+SPRkhVa7kGY z;zL>TYOEr430`)`2cMn!60k9jL>Y^2w{Z{~8v_@xi9WAi>;-*51ncIuSFjuLXX{o3 zZch7*FIzION(Y8Vv`SVA&JJcAY>@^h;X1uxOq~|sS_)hfFV+ZM(^EYSkhG6>)kX~6 zjVrw#GImq{nNM=%`8ic9cfIfzf%6>+tA=%}29P*ad->(%TC*qVi=$!dt*q4s(VACv z!?^HS?Pbe^z-_H)-Y4UV2GG|`#!ph}3$9P!%wC_v#J!g-jFEStY_a&SxMURQ7~FIC z?)7oConY!1+rH~R?&OA_ z38u`|@u(@klI1aDbH6OI$&&jtLZ&?^n7FtaFiNf=_>sLzgBpU@Y6w*sjpJ%A8F!V) zb4eb(O3t8ydQ^?3dJ3Ozo}M0~3~ScS(1eQLVxsPHhz{?kq^OO(F$*ts?LIzdNr6F} z4$(5^{3>eB(b6;P!STY-DTN?)st&Den>2_?B1;kfDHMvzUbjheDeMvrothb>GbMfWNO`P6N7+L^RS`nq>bk3U-tDZDnQmHw(yP=G)kl7kZs;|A(@ z2%CmHAuF|U%hs0a4Bw4i7%P zcPZ7ib9E3)$IvKt7$Vr#(ovWpsV;mzfI~qh7+*t>d#Yn;MMMwb8|~a<9UC|6^RS;t zUeQ6tX?PV@WN1)EaM>ENSIGbe;P6JYX|xm&E42G#*W?L)Z?R-bva3Tc22 z^a0cKD3YZMcF#fpAUahF@(l_ci&J%A0COBi&hMgZ-=bRXhaH<#CTdI!+1pc_qsSGS z@WI@6nK{;KjTzpz&=+?x2cJYK5CAjAAkr$Vagl$iN?^RS*Hu3NR{=%6jNAnJ{RUAe zN8XQv@0d|u?{1QPp8E#IcNfb0df~~GOemN83m9#SgYSqOBL8kNgv+2O&kusLeg7%t zmnU?<><=m5sX0a0!TF#H&7R?(yGcY`iG5AUFvYyYw=TA-dAa9v6JLzIu@Y|I!qfG* zeCIEX8|3`u9@3iY;(ZCtyl6yK#7}9(nB}mjJh$232l1gp<+_Y@^ff`JTV_87a=Sx($6JPcdashb zqW31Gs*7L)%D#xbiMW3Ew`%GJxOx7~s>R@FA#euwj?RGi;@;uYiqZLtUIc=2g=Ocw zzCxNgT1&pbRgyP-sBS1~`<8aImk+$--t$>`JF>(!sT zEYjypF6G2F6H0?hHMAPOc`t^U7KBAl&6BH$B{J-5_%x2@;2?$Ywzb7wqSjz;M$nqv zobs9;88?wH%K9l(KWR9(bK84GUfN>Hf)cQHj4tyX&8+Ha*;F~~U+7>x>^`zboMPJh zXQ(n;n z=M>}4DI5@yk~ZdR!I|yuy5W4CeU!8AL+p_ApcXUviIRJ5jU6vEPsaVFbMN5_YijuR zZt0ON)6c``_2oXtG<8hO6R~HlGF5d@@pD%WZrB}KObrcK&2PPCN6z_KBPpXz-%CZQ zJ}u$)9Y@b?kNOIe^C z-k*8c{=KTc#gwJ9i%+jC9ADuBO-)hMij*T(aq(-mF)ZK3dSutuKfa1v&2KM9TUTzz z-fN^bbT3aGqWiy{_wtLu-3eXjPsKNUHeVCot`&duP@gE!{Mto(DvHm1SuW*xoaIub zn=x9|D6YPy_x!n11tr44kThoM0;3NO9Xd2ms7;`F=xQ;SUnw_nXjkr@-Y6QVaxThu zS(GQNK}{VSqM|1m6Q?jPl=1&IPgukl{VwJOEMqwd#wXbo7n&AUlnXuJ>sP}5;(Ehxir3;etzxleZu~P zBtd@C^^%sY@ypBuZ~r23f%L;wq8)Vf6U454OKE(A!O;{Vl~>hjaa=gNYREpC8njoB zue!phk+LQ0$1$8UXR+SRApLb%>RZocrOnw4tIv0mr$rAN%??6qy*4Sm9^^I;KpA*Q z38C-v`%j2zJA$#cTXaKKDe4U}FugnNr>}iLxLMF7)N*s^Wk$=ypY)W1WfO*Xkd9yJ zjd`KV%9GnJNWP#EQzEz0N&U)SE-Pa;i(ZB>Y#ZWT%71ujas=rPMw|Ie+wXKK0-C%J zRdihR=>_cM!lBLePIK62oaRN?(Y%PK5Bj=bedSxS4^k_geO8ENxwgrC_;DXyl`Un20n9rX5gawC zEW!C#C+(XPj_jS&+auA9p%u?(ZChHo&Mkb5i3wd8I=_hS`NbqX^r`b*d(}9)(G#u3 z5M-*uqk$UV4Fi+nDw~x$ISZfz`^4WkpqZ(>#}~5v+PJSvm|~hyOr1dZEdP1d0NS*a zTnhRja%Ql{4pqM^>Q)_Q7#>5 zNxT5HB#i%7T+{wGVt=9cD-xH^))VJ?PnhY!6NNmhn1Lt3Wko19VUMC^L>+0{%q+fN zA++j~b|Ty5wLz`$x`mNss?j0hQbgG~0||fKG;_j3_*iaw?LE;+G0jSQG1D+&m8g3h zA6{P3<1H?-d=%@mbpL|^r9#eua>F^Vc2Z^y0-eK{8o|X@6+(8kf$TPn+o zp1$9N;p&Ok@C_2hQHWTdDimSAq(9$#KU&)W2;pG-Z?h38Gu5Fx@Sk#D0Dm+GE>#My z#{Kaudi+Wq4=(!C9{H}N0o*DtC@;fx`ls@wue+74M|v~m>i}DwfT&W`Qe#H#J;XLO zTO_!BBe>evITbQCQWLyqx-u)LnGqZ{rfhVEvQ+l&XU%E4=jECIkf8hY8aO)H#VQmVa!s9X&4WtqWe#}ZBM%(g{@6X{Eg(i0xp7` znqadp2n*q* za56u9*^d~3Cgq2#3`yHB`WE=_)sQCrGjvRAdyp}+&Q2cbcYUX$MM_O;t>G@6R%QEB z?v~^HnpM7Qh%gpPfX-x|Qm+!$EklCF00s1O@hkZ+#f5^+?@JjTZaK z>-t&<^qW|7c!8LWc?K?O8)cgu8KpJ{63&I7Ya?iHHj*P%q-#wEFgU)r3*#N^G7#*_tX z6(zGKFRz*Ed&&yMJ6+i~tAp*vOZpkT55GIsOs~XlqSEjA*2O*)OQHt4aRS%c$eEzYscin%{BiY<>MDK9QoKD<+(JNlYbx?=dA_gt7-4eJmnX^1JR|~jJX5-~WdzQ4Am$*93&^%36xNx}UbP2mY)7L7= zTD`My^WaByRq$i8)u)JMx3J$|Pg^5oG+LRO0G|Nj>~$X2M6KQh4tox0WSy*}ZY8u{ zu6tX>?YBbaz1(|;dj7wzi4k#`9$XWpShvEzGF8LgDCLFqPT9G9C^=}y-Gj-Bs&fXZ zCc&h?b(MziM}q~eloDg_e@ICxxz-kmBHR^@%)iF2#aHRW?Ka4ld&ac@2vklBpkB~{ z1ESl$?`kp-MlTyn>Z+R__r$7(Z0V@rKpJaMND~)$<)r*w=W>~1+vLRmz8XT}&dbYz zYd}6iN2=0hq?<&OnJ%w z*Xj?Va|Ti)<}S-ins_DY@EtD%H-b++#w$x{C5xQL$PVhwtOb~<$1520N~Z& zLpFLFYuu?}HO3Wo@*2P7v+f)yA*U`#x{FkxO1UPGQ_~A0%*{UG0aWuNeLJ_gM-4iA zGO^(QAy;)NqLX<(@RD+JrKth@%DrvHhmg(K|JEzj)r1zAJso8D|IS)Q+_5Of1>FNNz#+J_2X&-q#h$ajXhc;>UmZ|NUy!;S?!ua3nVmN>xFr{R+Ri?Y2ec^Dn zds79|aDLj(1<+tBYI%IjI-IYr@LD5_5buI~h3mJ4i$JfM2uwHljq38f&I(IJSOYvk zuh|hrXbASM)3Pqj*aJ=#Rfqj53M+vZgl0m-Bod#8NN+mP(VfxNP`z!)jj%KDJ{Wf1 zboo;>(cdEYX~i%5jDTj$eswA#oCI5q6)QFSp8PVAbQ!r%OiN>2wSq6Ye$wu(%w_qo zFkZ&%tl)QnW_Z;Ey67u-{pzn@DkGGHHbj|149-Rdk({}F#c*q%5WWiv)fn;ww<|HE*=fD=*?vc%GNmKV8ZWx-UBK0C$Az;SxNWJBZmvaj zspJ!7T%?~yMqQVx>_T$wtyW6yPD%I29gh@Rd{~N?H!70koSc1#u_*3TMA|0XRk&NK z>z2(Wxf>oL{8AqLF}%r--HT7L>F}-zLF}&U-YgBAzIW=4nDObuN!I8PN-!R2$knR{XSras-E@r}(|a=6sWdmHsvWu?7!|2*I-oDMyAUos!9toJe?HAq6 zj!!?OWpRV!xwum6>4W)%Ae<>I5P0u{R@)}izG zvjJ@C{HB|${K|5@h9rM~^7*|zDx*f4_6J0{p(PQ;+?nQ&UkI+Y*Zy$OM{~h17T`1P zQ|2Z#BGX^_XmpzW8Zp1$)VlQz7G7PmCt%Gj@-{ib`f!Iky`?bmt;@P4Y^tAa>H7(h zg`S_+ab+p8I&%}<(pM!EjpHs(NxOU-Sk~FzdJFwxdI~8@qB+e z)~bL+ik6W=TAn-!OMlj8A(l5eEXkrE#`=jszoEjAzle|XyRe+{h8`)B>y+h2!$#qT z{Gkc2so!^NlY5p2fKn}iRqopF$-GdpI~>RW2;p|RAEayz<>#^FuQ=ywB8!NW+(ayg zE1Izx%mT+jwx<6{!O!LK*r%4ZNtu%Zx6@yyqbL?UIE3XJNpvsOy$r z^V(m(B{UhO)qX#{~FUZx3w7^nB@YT%m2PcvQ05CJiD{K0<(*$NY592z&~a55P3 zylx{;O42b#E>*PG-;o8z}#7;;K%1z~J zY~AOn>?zm%skWeDiw-5LMue2rBXD~w*<|$ zBPee=FBD61K1kmUF2PVrq;TZg9wsyU1GnW*L~EcWS7X0T)J7o)O>+M)0(uf^9H&Dk z3YEw?vW_mvti_>ih+zFpuul<(qFyDvSG!1izS49>M|&JSik}~4#Q&$t1Uum!tff$mFWNdlSfZWd$doc9B|Lykq?wKPQFJYn}(W> KYRMg&m;VbAJf7(Q literal 0 HcmV?d00001 diff --git a/src/main.cpp b/src/main.cpp index ce9bb93..268179f 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 22; // feel free to change the size of array +const int SIZE = 1 << 21; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; diff --git a/stream_compaction/common.h b/stream_compaction/common.h index b2b34dd..e68c3d0 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -17,7 +17,7 @@ * Check for CUDA errors; print and exit if there was a problem. */ void checkCUDAErrorFn(const char *msg, const char *file = NULL, int line = -1); -const int blockSize = 128; +const int blockSize = 512; inline int ilog2(int x) { int lg = 0;