From be40f807fd0949e5d729882e382ee94008c439f3 Mon Sep 17 00:00:00 2001 From: HaoyuSui Date: Tue, 22 Sep 2020 17:51:21 +0800 Subject: [PATCH 1/6] part 1 --- stream_compaction/cpu.cu | 50 ++++++++++++++++++++++++++++++++++++++-- 1 file changed, 48 insertions(+), 2 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..80cf914 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -20,6 +20,11 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + odata[0] = 0; + for (int i = 1; i < n; ++i) + { + odata[i] = idata[i - 1] + odata[i - 1]; + } timer().endCpuTimer(); } @@ -31,8 +36,17 @@ namespace StreamCompaction { int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int oIndex = 0; + for(int i = 0; i < n; ++i) + { + if(idata[i] != 0) + { + odata[oIndex] = idata[i]; + oIndex ++; + } + } timer().endCpuTimer(); - return -1; + return oIndex; } /** @@ -43,8 +57,40 @@ namespace StreamCompaction { int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int *mappedArr = new int[n]; + int *scannedArr = new int[n]; + + + // Compute temporary array containing 1 and 0 + for(int i = 0; i < n; ++i) + { + if(idata[i] != 0) + { + mappedArr[i] = 1; + } + else + { + mappedArr[i] = 0; + } + } + + // Run exclusive scan on mapped array + scan(n, scannedArr, mappedArr); + + // Scatter + int oCount = 0; + for(int i = 0; i < n; ++i) + { + if(mappedArr[i] != 0) + { + int index = scannedArr[i]; + odata[index] = idata[i]; + oCount ++; + } + } + timer().endCpuTimer(); - return -1; + return oCount; } } } From cfd8346b57e0c9a64d226a8358a2b563e4262c1e Mon Sep 17 00:00:00 2001 From: HaoyuSui Date: Tue, 22 Sep 2020 19:18:59 +0800 Subject: [PATCH 2/6] part 2 --- stream_compaction/naive.cu | 58 +++++++++++++++++++++++++++++++++++++- 1 file changed, 57 insertions(+), 1 deletion(-) diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..1d12c05 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -3,6 +3,8 @@ #include "common.h" #include "naive.h" +#define blockSize 128 + namespace StreamCompaction { namespace Naive { using StreamCompaction::Common::PerformanceTimer; @@ -12,14 +14,68 @@ namespace StreamCompaction { return timer; } // TODO: __global__ + // This performs inclusive scan + __global__ void kernNaiveScan(int d, int n, int *odata, int* idata) + { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + + int d_offset = 1 << (d - 1); + + if(index >= d_offset) + { + odata[index] = idata[index - d_offset] + idata[index]; + } + else + { + 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 + int *dev_idata, *dev_odata; + + + //dim3 threadsPerBlock(blockSize); + dim3 blocksPerGrid((n + blockSize - 1) / blockSize); + + // CUDA memory management and error checking + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_idata failed!"); + + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed!"); + + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy from idata to dev_idata failed!"); + + timer().startGpuTimer(); + + for(int d = 1; d <= ilog2ceil(n); ++d) + { + kernNaiveScan<<>>(d, n, dev_odata, dev_idata); + + int *dev_temp = dev_idata; + dev_idata = dev_odata; + dev_odata = dev_temp; + } + timer().endGpuTimer(); + + // Right shift copy to achieve exclusive scan + odata[0] = 0; + cudaMemcpy(odata + 1, dev_idata, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy from dev_idata to odata + 1 failed!"); + + cudaFree(dev_idata); + cudaFree(dev_odata); } } } From 3b428a1c6d83d08d09bc1624f42ceb9004f1957f Mon Sep 17 00:00:00 2001 From: HaoyuSui Date: Tue, 22 Sep 2020 22:28:39 +0800 Subject: [PATCH 3/6] correct errors in part1 --- src/main.cpp | 2 +- stream_compaction/cpu.cu | 10 ++++++---- 2 files changed, 7 insertions(+), 5 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 896ac2b..90fcea1 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 << 24; // 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 80cf914..1beaf36 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -55,11 +55,9 @@ namespace StreamCompaction { * @returns the number of elements remaining after compaction. */ int compactWithScan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); - // TODO int *mappedArr = new int[n]; int *scannedArr = new int[n]; - + timer().startCpuTimer(); // Compute temporary array containing 1 and 0 for(int i = 0; i < n; ++i) @@ -75,7 +73,11 @@ namespace StreamCompaction { } // Run exclusive scan on mapped array - scan(n, scannedArr, mappedArr); + scannedArr[0] = 0; + for (int i = 1; i < n; ++i) + { + scannedArr[i] = mappedArr[i - 1] + scannedArr[i - 1]; + } // Scatter int oCount = 0; From e16490026ecf87b33b4a22b197ce321f7c62080e Mon Sep 17 00:00:00 2001 From: HaoyuSui Date: Wed, 23 Sep 2020 16:36:36 +0800 Subject: [PATCH 4/6] part 3.1 and 3.2 --- stream_compaction/common.cu | 17 +++- stream_compaction/efficient.cu | 146 ++++++++++++++++++++++++++++++++- 2 files changed, 159 insertions(+), 4 deletions(-) diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..336bfc4 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -24,6 +24,12 @@ namespace StreamCompaction { */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { // TODO + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + + bools[index] = idata[index] != 0 ? 1 : 0; } /** @@ -33,7 +39,16 @@ namespace StreamCompaction { __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { // TODO - } + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + if(bools[index] == 1) + { + int oIndex = indices[index]; + odata[oIndex] = idata[index]; + } + } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..25b63e7 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -3,6 +3,8 @@ #include "common.h" #include "efficient.h" +#define blockSize 128 + namespace StreamCompaction { namespace Efficient { using StreamCompaction::Common::PerformanceTimer; @@ -12,13 +14,83 @@ namespace StreamCompaction { return timer; } + __global__ void kernUpSweep(int n_pot, int* data, int d) + { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + int temp_d = 1 << (d + 1); + int k = index * temp_d; + + if (k >= n_pot) { + return; + } + + int power1 = 1 << (d + 1); + int power2 = 1 << d; + data[k + power1 - 1] += data[k + power2 - 1]; + } + + __global__ void kernDownSweep(int n_pot, int* data, int d) + { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + + int temp_d = 1 << (d + 1); + int k = index * temp_d; + + if (k >= n_pot) { + return; + } + + int power1 = 1 << (d + 1); + int power2 = 1 << d; + int t = data[k + power2 - 1]; + data[k + power2 - 1] = data[k + power1 - 1]; + data[k + power1 - 1] += t; + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + int *dev_data; + + // Get power of two length + int logValue = ilog2ceil(n); + int n_pot = 1 << logValue; + + // CUDA memory arrangement and error checking + cudaMalloc((void**)&dev_data, n_pot * sizeof(int)); + checkCUDAError("cudaMalloc dev_data failed!"); + + cudaMemset(dev_data, 0, n_pot * sizeof(int)); + checkCUDAError("cudaMemset dev_data failed!"); + + cudaMemcpy(dev_data, idata, n * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy from idata to dev_data failed!"); + timer().startGpuTimer(); - // TODO + + // Up-Sweep + for(int d = 0; d <= ilog2ceil(n) - 1; ++d) + { + dim3 blocksPerGrid((n_pot / pow(2, d + 1) + blockSize - 1) / blockSize); + kernUpSweep<<>>(n_pot, dev_data, d); + } + + // Down-Sweep + cudaMemset(dev_data + n_pot - 1, 0, sizeof(int)); + for(int d = ilog2ceil(n) - 1; d >=0; --d) + { + dim3 blocksPerGrid((n_pot / pow(2, d + 1) + blockSize - 1) / blockSize); + kernDownSweep<<>>(n_pot, dev_data, d); + } + timer().endGpuTimer(); + + cudaMemcpy(odata, dev_data, n * sizeof(int), cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy from dev_data to idata failed!"); + + cudaFree(dev_data); } /** @@ -31,10 +103,78 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { + + int* dev_idata; + int* dev_boolData; + int* dev_indices; + int* dev_odata; + + + int logValue = ilog2ceil(n); + int n_pot = 1 << logValue; + + cudaMalloc((void**)&dev_idata, n_pot * sizeof(int)); + checkCUDAError("cudaMalloc dev_idata failed!"); + + cudaMalloc((void**)&dev_boolData, n_pot * sizeof(int)); + checkCUDAError("cudaMalloc dev_boolData failed!"); + + cudaMalloc((void**)&dev_indices, n_pot * sizeof(int)); + checkCUDAError("cudaMalloc dev_indices failed!"); + + cudaMemset(dev_idata, 0, n_pot * sizeof(int)); + checkCUDAError("cudaMemset dev_idata failed!"); + + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + checkCUDAError("cudaMemcpy from idata to dev_idata failed!"); + timer().startGpuTimer(); - // TODO + + dim3 mapBlocksPerGrid((n_pot + blockSize - 1) / blockSize); + + // Compute temporary array containing 1 and 0 + Common::kernMapToBoolean<<>>(n_pot, dev_boolData, dev_idata); + + cudaMemcpy(dev_indices, dev_boolData, n_pot * sizeof(int), cudaMemcpyDeviceToDevice); + checkCUDAError("cudaMemcpy from dev_boolData to dev_indices failed!"); + + // Run exclusive scan on mapped array + // Up-Sweep + for(int d = 0; d <= ilog2ceil(n) - 1; ++d) + { + dim3 blocksPerGrid((n_pot / pow(2, d + 1) + blockSize - 1) / blockSize); + kernUpSweep<<>>(n_pot, dev_indices, d); + } + + // Down-Sweep + cudaMemset(dev_indices + n_pot - 1, 0, sizeof(int)); + for(int d = ilog2ceil(n) - 1; d >=0; --d) + { + dim3 blocksPerGrid((n_pot / pow(2, d + 1) + blockSize - 1) / blockSize); + kernDownSweep<<>>(n_pot, dev_indices, d); + } + + // Scatter + + int arrayCount = 0; + cudaMemcpy(&arrayCount, dev_indices + n_pot - 1, sizeof(int), cudaMemcpyDeviceToHost); + + cudaMalloc((void**)&dev_odata, arrayCount * sizeof(int)); + checkCUDAError("cudaMalloc dev_odata failed!"); + + Common::kernScatter<<>>(n_pot, dev_odata, dev_idata, dev_boolData, dev_indices); + timer().endGpuTimer(); - return -1; + + cudaMemcpy(odata, dev_odata, arrayCount * sizeof(int), cudaMemcpyDeviceToHost); + checkCUDAError("cudaMemcpy from dev_odata to odata failed!"); + + cudaFree(dev_idata); + cudaFree(dev_boolData); + cudaFree(dev_indices); + cudaFree(dev_odata); + + return arrayCount; } } } From a47401f25c24df322a45d9cd8ede0dba91ac45c8 Mon Sep 17 00:00:00 2001 From: HaoyuSui Date: Wed, 23 Sep 2020 16:49:05 +0800 Subject: [PATCH 5/6] part 4 --- stream_compaction/thrust.cu | 12 +++++++++--- 1 file changed, 9 insertions(+), 3 deletions(-) diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..c5ca0c6 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -18,11 +18,17 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + thrust::host_vector host_idata(idata, idata + n); + thrust::host_vector host_odata(n); + + thrust::device_vector dev_idata = host_idata; + thrust::device_vector dev_odata = host_odata; + 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::exclusive_scan(dev_idata.begin(), dev_idata.end(), dev_odata.begin()); timer().endGpuTimer(); + + thrust::copy(dev_odata.begin(), dev_odata.end(), odata); } } } From 85d47e13028c78e8f1ea148af2a188db9c0363ac Mon Sep 17 00:00:00 2001 From: HaoyuSui Date: Wed, 23 Sep 2020 19:13:25 +0800 Subject: [PATCH 6/6] add images and Update README --- README.md | 84 ++++++++++++++++++++++++++++++--- img/ArraySize.png | Bin 0 -> 19457 bytes img/BlockSize.png | Bin 0 -> 18482 bytes src/main.cpp | 2 +- stream_compaction/common.h | 2 + stream_compaction/efficient.cu | 2 - stream_compaction/naive.cu | 2 - 7 files changed, 81 insertions(+), 11 deletions(-) create mode 100644 img/ArraySize.png create mode 100644 img/BlockSize.png diff --git a/README.md b/README.md index 0e38ddb..c5180a8 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,84 @@ 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) +* Haoyu Sui + * [LinkedIn](http://linkedin.com/in/haoyu-sui-721284192) +* Tested on: Windows 10, i5-9600K @ 3.70GHz 16GB, RTX 2070 SUPER 8GB +* SM：7.5 -### (TODO: Your README) +### Features +* CPU Scan & Stream Compaction +* Naive GPU Scan Algorithm +* Work-Efficient GPU Scan & Stream Compaction +* Using Thrust's Implementation -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +### Analysis +**Roughly optimize for each implementations** + +In order to find the appropriate block size to make each method produce the best performance, I tried several different block sizes for each method: 32, 64, 128, 256, 512. The results can be seen in the table below. (Array size is 2^20) + +![](img/BlockSize.png) + +**Performance analysis** + +I chose a block size of 128, and then performed a different array size performance test for each method. It can be seen from the figure that when the array size is relatively large, the performance advantage of GPU is more obvious. Among the three different GPU methods, Thrust Scan and Efficient Scan have better performance than Naive Scan. + +![](img/ArraySize.png) + + +### Output in cmd +``` +**************** +** SCAN TESTS ** +**************** + [ 8 9 43 42 23 4 8 38 23 10 27 11 44 ... 45 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 3.9358ms (std::chrono Measured) + [ 0 8 17 60 102 125 129 137 175 198 208 235 246 ... 25659338 25659383 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 1.4142ms (std::chrono Measured) + [ 0 8 17 60 102 125 129 137 175 198 208 235 246 ... 25659264 25659282 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 0.485664ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== + elapsed time: 0.485152ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 0.219872ms (CUDA Measured) + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 0.19888ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 0.204576ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 0.214432ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 1 1 1 2 1 1 0 1 2 0 1 3 1 ... 0 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 2.5239ms (std::chrono Measured) + [ 1 1 1 2 1 1 1 2 1 3 1 2 3 ... 1 1 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 2.2133ms (std::chrono Measured) + [ 1 1 1 2 1 1 1 2 1 3 1 2 3 ... 1 1 ] + passed +==== cpu compact with scan ==== + elapsed time: 8.8561ms (std::chrono Measured) + [ 1 1 1 2 1 1 1 2 1 3 1 2 3 ... 1 1 ] + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 1.19734ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 0.7928ms (CUDA Measured) + passed +``` \ No newline at end of file diff --git a/img/ArraySize.png b/img/ArraySize.png new file mode 100644 index 0000000000000000000000000000000000000000..68842705dc415469b3b00d3ae5a5a8a222db12fe GIT binary patch literal 19457 zcmeFZcQl;e7dJYQB1n)3z6eH&1VN%k?=7MmqDPN1A$l7ni4uvLAbN>97|})_5fUv* zw87}Tgb|D~W4O=E_xHQ&{p+sxz3aXAkGt+#V%VN@&OW<-&VKgEJ6&zHtJJL2AQ0%P zhPtvo2t+Xi0-YmXz64Oj{nAH(k8_^-YKkC4|IKyamkW*xS_&XgWgN|k^+n)!DtC1g zPY{T;p{!u&XSp>UX?$a%@$8Hs*JtWGUE7bHA+$JVJpcLl($5E9 zK6`gOwi$b6`0Lzx8>5c0H@6(Fosagvd2{dc&p!`2Vk7rX4rXfE;0z)v=co4xqMNB3 zlNJ`%z6!AWrKM*|z|x=*Rs~kzV&JOeeCRozeZIkRC*SXZ2lq%Dy#v8L=XO+G@qYXy%!-@|Z9J&*4t`8X6ehu-uaLKNS zZgFjuO}^!p`uS^I>QXks>X9<%OD|l?R=zxh5-vIA8k~sySQjbz#&nv%x&^inZFLWm zWBN?wUK$82JAFiz^N2n_ENX;b-8hq~W|>(rsO={UzK}{%my#SnrjSgROU>o}RA ze=gvO$m^P6Ih2pp;H1;*L>Cih7Hz*ODA)JTvwLvolY4vdxxZC#rA^FmGjX|K;tlDT zy5r(iC7f>zHhucceAsYS@Dv+qot&G(w$KnNR~Upba-PM)mn8}29T<*i0+E_9v%ZYS z_lUbWwtOry#6gxfDF_2Z^Lm)r^K|0GY3$PKbpNc0~DHny>V9F`$miUH_bP%*JyEF*!R_XjpjF(J3iFt62{bK#%-N#)ID;S^_d#g$sH5a9X!po zvh9Q^)O)-ejinTn%Y)Lj`so*)c9}_9th3D*eLsyn&HMes;L!>qS znq(Eiq$PN9oNWyP)__X<$i@p#W#26hgyAmf3XalMVu7L zkSzieV!&mtj`NbMf%gE91BMsP17VJdQVC{bVLfn>a~UsjL+%U@!T()wqI{iM{A4gZ{N zTmBENr*$%e9%%QSokVWCV)3^{=3DLMkydgSXxN{Ar>W%_#f32h!62eqF?&>3JY|D> zqs!#@&|L2Rqg0h=hUM=ArTtoWW|4ESRwV2h&yq+k9_*J%LAZ0n(Q2o~xD8`=JgGi5 z8n^YP$G^#n&Dlbfz}gy2y{(!3P0j(-LHSyQK4FLQ8+Zk5!T9Io^=Fx4{KYPJ%7q*d zi{C>!TI%i9DSHa#`sc<&4pRDq&EHz0=5;}MhI4h{quTuK=Ju;pWw*A+cR0C z?t^MK*%eVP{7g&#sa#0yl9{D_$lKFn@0p5kM4=r$v`hezWe$#=PTtP zJ*5|a9@ql<@GpBj@w|!qH}Acw4Eg!wMNSXJ@F?ARzQ^injT{2j_sp1JlkY$D$a9g%-NsDWC;5!0godX+u}Q&Pw@&S9e#xvH5{%Gkm+gE_A{` zbala^I=0CMP>_e`3iD0H6Qx+;0nB{-@`AQf$3{rJ}KT|tPk$roZ8ff z-7m2rJbZqdO93gy?O?44mXT}Xht1AeI`yw)vL~C}xbR^2v&(!PC!(5a<|qn`#JxS>tP= zMYd9$?Znv0Or6X~nOV=URU>i>D;^vM(%YO>4S=EqFF57=mY^JsqT7yg;ZnNKs_Jdc zNlvgavEXTua7n(N!0FGwSdP9`kAQT69Gmm|8X~Rio2HmwHe2|yF=+biL=zNx|Mvf$ zjQl?&@;7T2*QZ1@G&D-eYRoaTz+XfZ^k{~RSBjo&obdTv0SGT9^O*8II}I~qv^DzZ zZ;1k=w5t0T*ZC3?(s)_WfBtbS0sz8b$LK)hRUB@}7A}z3|LfN$mG?}?kj_|k`{_pS z-BV&5k~2)@qf<_?5rsyzIF8`aax?LXM(@(*pyR}UpUt}k z7k8T(e>-uM=c_siL(`%s*4z^l6YaJpYssud3F17R)Gm;rNtEONIyCto1o6LuX#W3U z|HX#Na_Xyv8c&=f%qzdllYt`rKptCug}u7-6KXOz)Uv$sW}mTJT`TE>Ib$Ie2}HhF zARAP^x+x+cXW$9th4qDng&gPM!WcV8M@@f!2^GB?q%Rfo%3DtqB#npm@3L>Rele)M zEi7!1ChD?c+!A=0%r52ARqZ%HHmbb)ymVpm?NO5PCI0+y1OCcDf? z^2zxjIWT(rNF5Pqc--T8nXtj9Jf!ek&yC6H!&sT z$>v>uq31*%Pws%OWW|Qlog^63-&hj5*40sUmEr4SK6zjv9M=^~HC9zDHb*(?jjRSa#CboLKFt;9F3pFTX>^C0v`Q%DayIW0 zMzr&3R2FO}8me-OyhXp9ipVjpMgp7ym{Eyz0x2%ANZHoZWX$u9N1g5JV50v+4makM zeUg%Nky}|%)?*rZ>l#_^&Roz=35LAD3YnK?@x^2vUYmSo$ji^Pr14(Gf48t~3VBP9 z{OGNn9USF)^GTEhTc$Uj8JJj=v|iM=sPb@bWx6T%NhJ99zmo#`;AhA)6Z}=p4;z?QIKXX^xRD{UOy&U!mVFLcsHZy^)a=tO zlLgh~f{q4j8d7ivTbLfF-kyvQn^%IH>%&EbT7jo0{`Q2{Q=H)N!QaKs9wfS=RMvYF zgCHDYG0Btse2H7T?`r5ot<36W_n|3)Cfi$Z7tRP%BPf2b#R6sTghIOa;`i4-I_*tH z)|wb;!H$mp+725H`-~RSwey!2%AV{E$U#f%WP@-UNlwGJrUm-YE0$xNAt!nh+Ze>| z?rsk<4ZlMq93~Anrg*RRvk$L$H#P6if4ryT`u@zvf4$uuv+P2Hjzrn}Zz&!6Ot%E0 zYZ~f$dr_~jGcuU#!6qnR0)QEB3ldYCRf@Rz-V{N;qVl*cYINk;$*P=IdZ0;wMUb=U zSR@^%tIOIJT#NJclz=VlYk43OGFx+XB653cw{{)3SE9vN(L|@j-`ip#inF$;D=oBF8Hx;God%fSm)t;3 zx~909oaHKg7J$+3{%-x|l^(c=dm1#n5nr&bn@jhiaf&bBm)ojI${jzsUlpWa8mTtT zRR3VM#@ix!CN*f^XI<>DYG3feYz;}qVnn5lFLZt=pV_MY*ew^&!`zfD>beKMp?_=r zO_En!N=xWmOp^m0p=GjOXO9+K-iOQX^J|`c!%wo7(&0GM8u_QHSgYt-9`LckibEXN z0iAhp(!-#5I=$?lYwa%t_@;bD9ZXQNBMQ@MXTLYok&p0S3+>n4=M9FdrZCF|bkn^< zA@ZkY`6J=Je0GiMcls5sf%9`1`XA*}+&FRZtul*+z2@25ZtHB`2Pg~Qy+A+3uU z&~weCq!Rk}YyMm#)$grj4kW+)G%1(Le>6UsqtmEL@ye*T@D5bT5g;v+)2-PpYpu8s zIIJs;t}khk!Ti47CTB&jB1Sb=8hUmz$u9FgjJKSbQ%C{7`vxvJM3EOp^3=34S(E;| z?A3IZz$*Lyt|GE0|BmW+cWLY-6S(o;CAypN>ZSs}gcb0oy>d`KdW(OT34#~LbCkxV z(zon%i2}}Vwg1}dWmlvc$JiB^Bhi0PLkzK34i(D8ax zgWZh{)s(Vq31Botfi3j(6v!9>vdY@Dm_haXjaU}#vGk}Huh9sALLoWYnIS)7qwE|U zdS*`##9KFpi{k!F)@o+S1gic!u4po3KDi-(F`e<4E4wI7n-r|LfBj@G0B7^)a5O`??oO0IW{Bvl=rs(Lwdq&I-bcBZh2fc7gon19eDWRCagWkST<6&?ja385 zV^;r#S!jfwT#An}xtg}&j4QDM9JBi9EA$*IEot)|1!g)kPk>~u6))>W%!OY3FBveU z&-5~4ht0r~Oynms?_NX$arNtiStRHxIixQ<`Ad|`-?|Ru)V!tvDBs`?+42xUb*G=G zeh(j5Rq@L&k}7wh^^dSS5Qe-a#DlK@iIjsv>)qk|B{h~oT^H8P24*IWa{iJnzLbKoFlk~q;$D9zvzd4WRq#?xn}4@f34ff2$`i^!64d6Wqwf5 zD}nVc=~ul-ib}GP*^K2+GXrebo#sQ&*#MOoV9)tLE>EYQ4zkYr6axQRF8_zj0R-Q2Q`|L_;Rk8x-giEPt0~ z$j7z@xtWC>DUw*eif|jFV#s4=HV2kCA#3WA=;m|4DPe}6FD;N0ugGTrzCNKLfuZ2T zN&p4?B@2PIqbQnM!xA>~cCBoX9K*=VQ zgN9_Sw*(j%30N4BWdG#|0ll~Y)C7Q1aOo?len9p%nZ8e)no&7cUG=aC0;KErR0FNR7jHCi(6INXF%@2+>M0NfcWtWrdA7W zr2HmARUkN{x%tDNTRlZOz{%{~3%#pD_{7*xk5Dwya8hUuIb99GL$YOIb^}>5yZh?~ z~-8BfPDm2*YJ2z*v#|V?SK^Anr1iZt^9|b)( z$CL?%@hj=WeE)p?Axipd1#n<^QW_4VszXauT{&oKW3@fv8c7lKO?ZS9kpC85Zz3@rEtl~2}P1&UrU(1_U%f;6PM*ahW_Yl#a8v064vJt7q2}I{?Ay{7c4qMn8?>NC})lf{NWp zq=hLYjV~#kti?mm4bqb&*HL*72x-8v`M>rqa)BiESFU8FrX%OxjRd{8PaZtd_}xk? zO4~GG93ZKaNrRcp*n&}VU~d5e`#F(2v0h?eVSxyuerhgVFCyeGv10*9d>Ds_dlE$VAVA?3$ zQaSV@eJNS?Q9z~wk`4ewq37s($)wQHqTokdzzLcQ7(7O{L_1PM>J9HN-?nTfyf8c0 z(O68mKJ|zvG1)~8PD()Ee&5@COZEFb4FcHCUvRVnKH4t7P@PVg?|C8_c5UqU`vPfj=H0R z!;@-4-2VOb=TaU$s8!ej3f^OaPmGc`S1(iNi55q^|t&z+8l@ z(ds}05Cl0uyZ~Pm=S-(Br;#VuqXy+C1*KdS_upZ30O#-{o!)cK1p)Zy?V7b2NsE3Y zz2#(Wp}a@{pokWvMFJ?H0sHfd_newK02mYUkt?*2gudlvSgVddQ2hn~7$8C|a;ONX zns)_~5R{3%iXW7Aon&~pIRFGYwLBcSuuiBrt1164IVa?Rhau0A{XQV#G*;4+ddLM* zoWbzRxB^zILNesMwf^^qL*Ffj{vv>y48@CRlD?Jo21GlxqAc~@rvayelS$pp?ZyC% zY|i@zkb20Zwox^~f*%yCiL{`1B&;ca!>jMXL#_Z@y%NBYM@VK;{DoYInzqF~0t9xA z6h4aIe3u%~idspe1(OkP@f+>lJ*_BNKOB&03}9~MmBo*a9xe;!4YD?HY10&2v)?z!+f$e=W8hXE=x?$y3+=LCf28$t z;Ms$?y6q&kFNM8d`Uro7qsRwEcx65LjG1y1IQ{uC+H)>;^TAdlzRM2Z~GJ z&b+T6&bIHYe$-$39bvJ%9%t-fyZ?ZP_^KiE`~s+&dWJ8%Jvsy80t2ZT%F1rhh2}v$}!`;37=@qbb z>1(I=U4%&Tdb_hYVEi=d)IlW-s$g)~`1RS#QjqzS=ckNgAql(Zvbb6HzKPYnEyL#= z4;6wl8ejAC$|h= z$nJ$$ESaYV>koyLO=UHz2hTQIoi;<9S1HKDhwSacsD7uv53c5qQo~lvCI@PYcO``_ zF1xArNx(5gdmRgyZfo4Xh{nj zMcN;8D|0p=yO$lOdHg=bM^)aHlx4144kapaD`&R?bGf3FM8^aL?7>7C-pF#15rV{3Wd3q#~VOe41l@E$M^KeX_Tq};Q8!ep#`dCj;B%I z7U)h+{PfoY0c{1AL+2Hl4_R3X`G-kQcBG^zFZDl4z?i?89U=3oF`KCrxkl9_FrM78 zb2`S}P-pYeoIO!s6{F(r(!ge4lodSBc<-}&Ed|@LZ60=UvW9pLmwq*@pejw;=SIqEMycf{? zQA90u@2%lqjiqUZgd$^=Z4QQpN`s}m9?^TN@0WD0Njx^S^M!Yfad_Bx-ZROP)2q@M zNV!U>)`KOeji-WVNP&`Sd7??kQ>%t{h*S3GnfUa&SvwKsA2;=nuvC>Rwn|bXA;>Hq}zQW&(4a@^<$1h6c5Za}N@PX;fmh8IJuwbY->Ol7J~Ay%;A3E$uA$ z5^Ak9f20X6wQme8hr=md(z?q>t-1k|~y`Y7(&84x`}-wt_hJz~fUIXRAozN1@qo9V_%Q2w%KN|-U- z-uTM1(zq%YE0!QAczmR~WM9(Nev2{m@V)uWeGi|?u1c0${)@^7w2a+bd{wpST|a4W zO_W1GF9oVrG!r1%y>doGiJ%0=JA zE7+a6Glml0xoR4ntN%j%4XRRQS`0b~kD^B9QMiZ0hjX3zt(vS zd~hiygov(_=0{a!U<<%Lk}~%AmgkcCJQh-5`nx+H9`xy?_&cYKsugwV5F&e~SL<9m z|6@7xW7T!M%A6t!l16KDrPv=g|CTS~<46V*5*8G6@~gjq1BSo9dyw0O{c)?{_?B_q z`;6Wj5~9XSUunPHR@skSMIO>mOqpDqyR#~E8FtA}OW2d?S>$tzW%$gcjH|L0g;I#u zI|F#vuhie!AoP|6iaUyop^l|#b_18_u3oO0fyZ2#{gkiheoq54qr$zkAKxtnxop%X zc3V&@nfYn5cTuoDmtO;QTl!KxbsIA_JAP;tz5Ii4Pai?ZqwSxOxK;FouBTJ7pRzKX z6ge(W3&o|;w*Ut%%#i!(V?Dd=ALi+1y+ldd66uMgd-lFELH-&l;Y6_cyROVHyr9)+ zbMxCCrXln%tBBw!a6*P91bwb_5MFUUc@*Lp2E(~Q?BQEYOzjCQZ1_3FCTsW>m*vi|xT) z2nW9=&5-?8i9{7{N@&73v0QP-vcc~#{{+7emo^BwJR*AGJo%Nb*3AC-6yV)ZRn<4- zNStx_kCn8g$tlxqqkxEpgX00rLGlYu%P46B>oWwv_1#&@`NH=n(nfYz`apJf-vq79 zjLp#Mw9=i@it}TRVWkvbD^3P3h0JGM9YbGAuWqyI6<~ZVzFFP2X=-T6RyJ^{sDzd2 z*7z7=vP+lowc5*d@RvH+?M2T-`=Nf6kmoEo^>FX7YN_jviKj<0b(^NvW!V7RI&`DU(K7u~@L1gKuJMvTuUUW@EJL>Cc8MidK|m-(BZep8KX&x4|vp zP8@J>kXiU#FPs`eFQ8i?6OT)BHL&(8r*6Y*U-LA)oCT-Y@nig|7Jda;%d`4)XQyw= zZhx#_$S2S@^%F*|5X*kwlP)%~K#jY{l<{@l4gQ2_-W}$t@w=;rmr=)omD0hmIQkFV zq(uKBAhz%N1TM27Sls3KDMMZzP>PdfnOl@fNZDvmV^)62dA*@=w1M`h!<=XKk?E=V zn88f{OTmglWFq=wYckpZmFO#s`cx%h)g{Ly@=>SvJKzszhUbp9-~-Dg)NQ^H5Gr@$ zdz#?#C+aqTxQ~Nq%2nPA6^2}XDC#z0U@Q5Gf7Y&++xh8|ENPj(KXNf+^L6J-EVta! zL9$cerLK(gx9$GU|E=0Mx-wy2qwd4(sBwzB$-~=MaixuOpIr!K_#VvO8SHO*%a|jx z)>i&DVdv?^_?xx195q*?*zf_S#O3RFNqyKQJ<8rH%3v z*Iz(`fB&Gr+^BY_A^i=iU?Nqe|L)bIl~i56dm&yFyOpW^#9&;C<;~QkptY^R3#$i? z{n?9_6LV_3m(Lmt zQ_b}h>eVmK-F87G%DAAC`AzuklBmDEHw2yPjoLBV;A-3Ot~-N&2D$21oNb?kJbt<4 zuu6N4!6Tr1M22T~@iez;Zt$^*ziXN+ulfBZfwJV=Vi&AnUX~oihimLy|Gc38O*4cR zx+8j5Mc*q%ul2r*d^{q3YA;nN@XX!hliW;($()ho3&Rb%Ml0lD;+H0TZz^;)`lhk9 zTBHzr*Z#SdKLrdm-uH)nL4$Q|Lt%fGq4z$e${Q_Qs8UVA2AnTh(kkF4eA(+7EHPm3;-sG5tou+s~xy6WX5(+VD%!ZNT+@BoOwu2?Yk#NKD zVRQ>z*I_#}epLUe+2fRRYxKo(^o6}sEc1N0XCaTDCQTg`vy>7MVdKtyWciOAJwHCuYS{z%FOl8z^KI>ChTXz>~|5Ul{dH7 zor~pE%ITQ!$%+O(MV|>S8^ZU8eL{+Uwtg&u@5}LNV}u7syl~H~T*PhN;Ky_C+=}Ji zL7lPBQ_FukV%hI_j_lkU6?4>q|AGwn?IoqDMNo^dI0w@h@=r?1Gr^Z1=x@FErHjAo zTyJL_{Q$CA{j^rlxK_U$O%37mJ@^Q5{#+8td!c-+0BK)7iSl+PRE;1J265Gp$EFo{ zYnZm+L^0m(3+qdLVyy}ey|UHsp1xsbwOqXO=?gw2GCagjuKu}g;?3a# z3UuXF=yE-^#l6umLOusvvYg^GLU_|mBbVAboc9BxecdO|O~dtEYI|nG`%Uqq<^pMW z(`bl=J|q#BWNt}FVzJb-UW^V1(z;qvGVi^oGi24HB^qpS)f;^LV@!o{WsFlR<&AM& z*M{xzT{k6TPxe(o-nsQvT#YznT6YnjU0YF#wla>ZDCsmFwS`4sqge&@;xtE2aMPgu zx^TwQ0~^tiHZ}A`2;Rp8Pjdr+yppA3I#z;kEZALDW^7i-N3!?0Xz%03Un&dd^nJ?G zeP+R%2Hu8ADAazG0Pfjrq!HR05^1k7l9v-p5kPtK{!Z(O-J)TbpU=O>C}^flW7gOdf}uo)S@^LTjTEMNAPO7SK12 zIf|#aikds}X&U>`n3P7i7y`O1{LOj2*Kmt(bdho6uLlIbO@&XJxV!P!M$|n2Nl3hF z80&y8;Wfg46H)?Ox`*?NmHh%gV1pTXBPSyrD%_v~!k=?Fbc})4WgQ4%{07diszfXi z@pE|8UEjy-SV1N$heJme*Y)FCjoYyh!hMSk zx_5fc=5idl5w}P80*YPg1EkGmd3!7EcSzv1ZemZ~jt6x+AGa+q;Cz0{=i)-$~(X?G6*t51kwzRT9RRBaWI%Lv`G zbLv}Cy~o;?lWrQ5A^u;=EZ#YT{~4$**n=)U!Wahm8W=YBWfu6Dghp@t!0{{NT3k?^@4AG4-%hU z5wU0ZEgNIRJpgH6ml@$4oK62o%h@yAy^t|46d0sqk};;89J>kT8*b7#XjVvU@`vAS zs<#lBmb=?3)5;?#Cw=?E`QoGinjuk(83iD}oc!6YTIjMf2?@w5>(Tu9-gENFkB!#e zRwI}|HgL8%EpfrEt|Y59p5}6d7`C@Qg}p=^^BD|py1m2aY_zbVaoRG$b&aVFT8XEu zqiE-oT2|ChZDDGIwrW&A;+c}-#XqR<_r-L=UslR$Wii&=;TR+ z8q$r}5_QTBAp{wS@95&}$d0(^i1<3X&Vkp05PkwLJX6uC*i1(J5?uZX7hJ~o2rh3# zS1fJp3KrOfBPy15s{{-DCX78^mGI(@z9&{JH^V(OQHj;g&+3o(B&yf=BwE&7Bx-ey zJtayuj$T(RkH22=c<`m+##-8}d2F%uY4eK53riXUf6w&Z0cSsVgaKi-<`Yb=_F?T| zQF4`9%;qG{R&$~xLdvl&(H!R&Sz=TLOXeL#j*c;jB3!EBG1aC$KU2%?-vl%V7b31T?FTI^!M>kM*QN?Fv{Cy0t;7Y$k<2Q?%UpHAfcBq;+J*M2%&#N+UcCi6W)tEE{+|nbTKBY^mkW$pO-y zK5P)EeME_v(4v}5E28T7X~cwh-1HP*stdlyOO$v;bb!T~LwMAETx;VeWg_uF!|003y;P--cRZt5X*CidZ!p~( z3eJy18VX}w#Q zm^r01=4*_D>WKa+=%H5H~F4dfAdrTI5o}Ad0>ej+*`*2_VWTA21P6uANaqF*e1&`vz=MwoGoJQ>@+Dc36tGW6)8*Ac9%g<7BJ(_zT zO3ZlIXKaW)rT(baGqHU;|Gxj*jLk>-{UK8N(Sb0*lz_P8MA&?PHhjz{983-V{BH|2(--@4Nif2QvV9cxQ{AEBQB~F9 zGZTcfQ^l*Qn!^hwQ=WMb^lNig>EPU{J%ghziNCN(0(*fvM!@%{T1KKOEAaU#S%ak_ z6deB)(q1CeITlA8`8q3dT!|G)opx)4I|-GxvKIc*ZD$cOkjX0H7wf2 zDd$78(Fdsx&PX_?J)BjW{pib6JyB8bSl0f8?#7YDeFeF0JIH#8Tz9^LXC~{P5xO_4 zscDau(Ln|b5e{7!q;9&oOntp5^s*jqO=DjcAnkLc$nwScPn3$oV~3+#8f=x%RMSoG zd{KkXcir7>LFj46u+|!m3bB+#+%}9k2~Pd+Pm6hy@XeRb;a4<0os0dO^&?_Uc$9)f z#SE?r+wY7+gu09u_uXsrW6k)_LkePz+yA=Q&AqJ^dbzJ?k|LpRw5|COk{;33+|6*; z-X?kfjLU3PQut319Gj0D5{I?yYs`XMytQV)Ekoi@2D;Pp(RQL4jV1w4^$H?pwWn`& z4;z+P9|CY0vtRC0uhrN;{u18%;pVYJA>2eNf3NK3dfBApz|_j8G}PV){vi!W=X!en zVol7D6JAx+96K`Ny8hTc)9Z8EUG~DVq=S#>;#~MXWJp%VU{c4lgrmZ+opH!|NXSXt z>0M2B{j*7(XLs1SAH>Hl3g@@i9}m6yW4HL|h1E2b52jbmO*#fg4 zG76omyjW5NEFyFg5$*S$psg1RHV0n~LA#!}mH9lvoiHP3l6ja^8XPtBnM!zEXe)+%c( zROQ>1cU8(R70w7ZdO68eD!!^dmVDEs6&mSkRO?OfnLt9#y0Sm|@xylKN!=NWvTAx) zEf~0;#*0BG-v-6zv3}X-^xaz>sN<=ArNX1mSKI6oNjszzai!9>xV+gV&VDlbd?L!P z?kUFj7FISH<2V;i%U@7X;B%!GSCXf=!j9r@M>6%)XENvS^m`JnG}i|nB2ErN&g$6D zv!Z{-uxYuUvrz$BLeiecBksNyu>oGBRlWts0Pm4if!D-*7o0kP_lN%{RVm(RKYT&h z;n!DN`9)qH9xnCU4r}1@Lk%#~8Z3c$t@2LtLANT+(9|jEX3F0WKt1Gt`}I^{#{WCF zTK_*dOj$;KHC7{`##~DOD(OWLXu;$^cQ+C~-5?XvP^9ZZSGwzevH-6(L3b{ZFIj-3 zcR2Drhkx1M3`_BNOFfjDnmSbNsMESJ)6!IC)oN*p3eq5VU8?9^8~sl=&{YprQ_A%M z%@q7-rETX@dn6tC%7@+!NNF@DuUjUS_|3(Qu2U1VS(x&GgSWn3KY-9T@IGi)arwhexugNAx;cnb}W$Kkd$D zD{0Y<#1f85gAXlg`wC^;r8~S3L%{ zx~D1U`=L7fIAX=bd$0?~e#JGN;4jMUbkY|20)B946EUr3aZXjv>bPl${cG5Z9b)Vr z5mAQETp!RbI$yiF^!l_t9wrn_pTGEywnr?~ko^Xn;{B|N94rT12Hw~|bvTqclr22e zLWg~8v2s>WNUIKuN@fif#2;OBY9%HSul>$HyjP1|7w$n2f1O$3M9^Mbz+KLoDEkZR zk8?x5Iv)Ya_DF2`Qsa(-)^sg)vahoa*P1Q2H#^n zIZ*3Y{eqYxd@PU8ZA@GD`i6wo9-qG!LvEz~pS|YZ(W0;6c91HPuJ7D=4s?~&bik3M z?Zlles%iWPc@ChBSOBDdLa2l z;LU$1Gya!59itOWiId$Yy+E^(dOSOIaPz{p7wiCE?T{F(6la{N=VQqhP+yPhZabmbr^1S$u^YK87&gohxN5{56l`{fDQ z?6j9h!yLbyU$ZyLg!)7}M4qug@BJbquNto9(1N0F0`E1Qd5(6eEW{q5Z{&NcBd%0x z=OC_pHe!6tY!WfL;>cD66h&dvs9TT@wZClYz z#d9~eB{$$;_Fy91vb3;d{)0Fu#{t++@l_bpW9327T%$U)pV)YWJ;UdZcPC(~rc}+! zx8prEBC$9~?<>pIknSi0lFEXXt=Zf@u4n1F+-i1`@u zW#!>*;xco5PLK^%e^4$2{k_rC1O(mrQEtwTXLCZbQ#yV;rfN-(kn9r3*UEb^U(pI#Vf%ij* zdsc;8)14j@!SmwlHMgwBGLJ`VMk5&VEGYTegfbp(U0Ba~KHdHfv=WMy1w)vxPtOm! zu~aV3?8fMiHy26IRH@0foDJ6yOaBNS!#WjsPHvfO+m(qyYxP*7PZSmR6+L1Wa%Pt| zj1sPHsYIjSpWu3&2*wRvThffh5JF|Dwi*A0t#pBYS*t$Ix3~xH$H5D8#2$Xj+fHV! z3AS}UrL#)*g<5oRvy->KK$?P4$dj2j0klI`iNWXJk#_U7ZYj&l(gdtCD@lovugh*p z((a8L+5P49T$lF7NizX_!%H46a7UHCN^8!GQ+~_k1cRH97M;1>u;J=|gEl(kro`=V ztHV!pb=XO!kPw-myeur2&H>QIlkf;;haPqfl$Wc|u9%`V{=AE|okdDr`-Qc-?rIkQ z57yV;G*8h>Dq&V|A_BxY8|}IUSM)(bqn-|;))~G+Z>u$)iK>R1?FB)sj#}hUvIvK? zo@_=bEsr?(ano!_kvVMJHal+gaC6G-BJTs=HuK#^0_<=O`t@$>wh5h}9Y8l>NpFugu-WJqKh*thQl`l(IP1HC9TCk z9oN~3b&3?MAab@sf-*wGdYwy(DJ?j-KWyITflwP(p%_Io&P^;M$n0)@nd?Ml)3w?+ z)UPfF$DJh42gGSL%}QdN;QrV!W#xF)FM+$rr;-A4Hl5%3_|EJQnX#V<-#8Ox{yBzx z2IZ-Ed_@kg2~no8jT?vqBN7@7E{FH zV=WhaHVl(;2v#)Z1 zHZtQ@N#a?f`h1GapBtl6n5wvpQ@JzrCa?4GfcCuL{R?Z*s?E)}pr-BZt;^%wjvW0a z(AhxiNoQ(X1=fn$yLGKylNE1VloB&@E=jZ~j|E*j!F%)nxo`@n5!9m_8r5F>(74K| z7ca(&C8mAGXsW(^i8XE)ul-m9-Jr|9yhaC(9nZKS68Ek_X6wf3_ZCXiQO`CayR5pI zgC$%m1!ggRIQq5sYsGIOhf6_LiK7QL8%8LZtwGS=fMb6RDAqffZw1FO9tr`b%;cOJ^EWuDH27*6Gd8Ph zd}|sI^l%vAM^h))D#MfdL;2DGfye4n!;@RP#Eg1;FP2V+APJcDQHvnrlp2=>V4%4= zf#6}oy_xeszhlhjLC0w-9^gI(JIz*CU#F1YFpCtE#40!^cme~b?mwmcqWeqXJ8p7^ z()mnSCRpLSFGyrr3e-?k$0;pD`9<%ok=k5;w~D_Cw8~I#)mZmi@ZoVG4hJo! zUkYq-!w?JqVrP4@lQzD8x*D>*TH0|y5cR}@8%}cYvybq{sf0@tm0gBJT73)~68lH& z82|fgi1gYn?CTH~RyEMDvRe0`sGE8%5ItacMEr)F$hI189c?+3r?{nlo7)xW2HrRG z)t805MalkMj*MGDa8OprRPh(b%zDyx0WEk-! zY11XGwI=z0nl43Un7FD9ex`o8)YHMHKF_q@tGD0M z4A9SeQ|9z}SC~@tE{pqlJ9HzLWXzn)u<^R7cYuIiG-l2AQDOoNDD{}AR;LZLx^+@(hVXd9nvW+-9yfxpmf*JgG$#( z4GzOF_dEVx-L>9&Z>@XZz4!g$8o%coXP;f?v-jDXC@l?T${P$fh=_e*KbF`$X2K1B9kwO8`_ACH?({Q76D?9j zUBy0Aqw7YBR!B-LT%kXY$cRO5?#z__mZjB77Z&a*WjbTJye+4_eej%l)1;)bz0z?c zaHbkOgU<)1C%W}0>?#ovi(8l$aG@Aa1FYnI|1HoJA|e_I5En4Urc4IVDYr%jtc&R5 z1J>KXPtyP3m-_Fkk}N$sTl&;~@?hq+f7yv}EtbzspmBL6X;OIP0@}P}yw@PGV|%$g z_9W2YXf|-NlXz_o-A3mHK1oUPU*6TPKc#EPy^57Tk~9|jDhrb`iS&YhSmf0IFy)Y( zB;5mDHnp>q)f~xAU%lX{T`YYqy`Go%Xr$xH+6@qwTLh)$eLW!wdbb5ViH;k-a)=ir zPsOf3=8L2-4cSk#T-?+sUMPjrWIP)M2Xkw%C4=^t!^RWe?QQVNz?{enAb6D}EYy|jpNkD=AxbtlzRCUT^Pm|d2Im$}BhbT9nD9!=J5$;QWF3m8i3jZC~W zQDQY&LH~X4J0Mfydye6+oRp$W2f@FLk2{l1F!R2D)E~7ad+~|a(~5PGA5XgI`txjx zu!?9|^v%k5U(AgWN3sXIkPIcV;Gp)&>@gRXlOx&r;Mo*ZThI2*d-vO^LKnq69W5N~ zEn$z%AI6h*thufTg%Nw!hKpRnR{Vub(bv82gXqrib%A2*{veRu9?W9u^-4pT9h1hY zd0Eu%6SEIPD;r{@xzs|n%zT+)EI_4?&DVZ|1$CRS&z;f{n^8gy)GenSL4CC&8POFC4zEpz`Wz!9mlOBtWHiF#G*|kj`2FumF&*cP z|H22k7*|heo(E~Z38q@Mk%6B+2>liO!fKr7AQ$rsFsD6 z=N)tHk^7xYlyo`IQ*3UHc0#b&lh|S^qC@+~x_g4$Jm25he|Yh-pFGH*-;X@Nn`4NE zW3OY9d{M?ejaM+N30}6p9flJAcILAvYL$JCD06uzw5TzN&71R6F^RlH4D5c39M(6<2&s?hdIg{xuaH+)GS^e->On^=E@uqAHv? zT6gk#oE$L&+LRSZXGop-L?PqN8+-8F5eL6Pcjw&CD3rawHAL*!X9EKMm-?8w& ztx)jMFKp^rZiuh*ve`;&)8gP=^z+QB9vYRYC<>Cs)#oLWG8a>-n7oYH-*si=h3EJV za({?X;1n6|GVJ&BUa&qcz_2YFb9MY-nATdX6lFm{{t#Cf^mX~-jb~yqp+u`IpKi}; z>!$pi4A6A^@_I$GwTLFm>wb50v-5}F!em+n7Tm6v3watFtWB_QJ=g;=*}IDzUecUT zGi=rscmke0?|;J-L2g=`x9az&^bT2p#n8^;=7LE=NZ?3ym&fr!T9~zpOoUEw zU#_z|5mOvZbdIYntB7G{uTQF0C+(S8%dP96zr1jd5Kc6%09BuAe~yY$r4tDVVG%>^ zQxs(KDSQdrzTWD3I2KurSaoA3z0sr23Ey+pSevzervH@N}ZF+amKfXNR|33p4LwC#pgdZuVPiR?7`&NmwyE}5xqxqsF>Q8@B67} zlv1>>vwQZnEiD3{no?iM)qQ~G9~&Ii)4Y9Em3+$cef1<1c~_s+5Bcd{%?0YPH2Y#% zC0~&J9+$R?jvy(|`gdQQZ>MkBuZPB_v4+n{-sF$WHha?%KCJx9ZQlB@t1pgPjW=|= z*w0yJ(_Fk9oix|fGb>+;EfROeKmT6xFdZPafcuj*brM?uJBoCr>DMk7bR}SyWk&4A z{H&KGH~i7QqzCfWQi9&Ei%iz`xTU7f{Z#(d-r2yN|J`VedG~bLt)E5deXcJ|{dAUR z7|uj-|B-3p4GjSfSIN?1GU~&#urP6h`5!A&nPcRX`YTbpKgg{LUYzb9`kgq zC<OvIg+ zKevp}25a?uZ_~e&4@oP~UW*`rfM6vW&~H|EJMHe1-X9T&c>a_R7T*C}@IK(Ky{erX z8RaAm!-_EhAp|5#vo_g%xqbWgV20lr13;oLFE8jze32{w8RX7xy(XK1=l`tE***^u z6A|f;5rE(0lJ+}Etf=5=3&044g9@~`ZogVd3qbtfD=6!0o_I`7*4yeG`@Y&1>mARa%JfBCq>Yu~LFm_0Fu&5?Yc zivVKYFa0k7EC31rlc3Rm^Gqh_u<>(2ZNsO!(Mv@KSCp*)%j z94M05MRk&9q0x3sZhN9AQrEy>2y?v2efys%j~Yx`-+e17xcIT|11v(3ADDaZ;P1_Y zL{0(%U4!r0laJ64;3do1eKkm-{G>|K z36IBQL4#D7fl0(~fg1bA^9QXld7sco+aNf1=w|v{@HK!GMvxsn=ACb9?6R&c9`^aX z?gSeYBVyVLF(=%m1MW)r9cJAGoA8HjroHSc2d^<(A_j}-#9xin#~d&LtHXrQWQ9YP z9g=j?A^pEUiOsVt_;-`cd#%PYQUGLO0kW#v*1a5KwnBd+2BMeymhTZX_XXL=m1|`# zE%Hr>IZvUkqdG>w*b+O)3MAf+=e%W9#0x@l)s|Pb2-dpB0kWg=b$^|CS5f_bNBIgR zL2}Yz;}3*8zch$`41)S4{CgV=yuI~XufsTWi*%@R`JaWzla0K#6q{`!nP|D2+Cc#J zZ2-YNrS#%Zk_m&&LpmTdVDY8_x#!3&cDT=Y=sNeVB_xx!<$K%c55p;4V$wXvtZTd)@?Gk}W zdvatWmSgHjNaz>z9n<#R48khMPss{=Yu@PBQx?K{t3S|eKYaC%U_3};c4m89G%F7To4srUCS2XFEgJ3(K z9A^6W?B=5|FLyA*SU(`|c^1hx3u0{(OTP&0@puU0vN}6dowV&HVN}e|YB16;ZTtD} z4i$*2s^@OaP!U0a&i@w2O};yiW(0I_{E)1`&J)wi=X;JF5&MYFtTTH>mFsB| znaw67bz4N;(n0F5C@ilQS`ceLYoGN`y2>iTATT;-5xRBl_w0bCsPYm)gW@NAS#=cv zZ%ij_+4MuWzC>6etoJha2CIl@(ye{)^k)o5&1jfBPlN1P}r%e4IH)IPMvWtzF znOpx`oHl6{Hf~30YGMJ1CJKDkMo>cUvKWOg2wH}Y>$Q9Z&cv@lb|)`tzCb7o&1*D~ zrhuSW=u877WGoE3Z5bFCifza8LMDrpsmp6>(vCOBbl9HALdu;f3z5!qcg`RFXa8W* z;Wueroz%MAS!4VZa7DP>TGyWJv^k!mQ(3{z$jIpP65!9o!Y}zK?iT%b&p5 zQwvap4eCbtq#!%e^DUX^1x`AUFQ>VP;&-j6Pwz_4qpXIN*yK>);W$jmKLKMsAY}X3 z!^Ue2bluhb;*Nt{f&i2KrxEU9T1!c%SdD>D{*UmdfgS(M@;->`zz~ykJn7#}uoulm z9$BQ57eRuPkHUV{h~&Z85;T_VLbpDB5+9gQMf|guX|jS=N86f}l(JGMsQgL2N_O$V z+MR#)0tDtvzx3T>PV(DlNG4?H#9ji%`0vhr;Lc9?{Y60dtBalurI*G3^jVQiDU$KI zcNt~jSpt}sf|vTA2wcOoJfd6frZ2KA45iphQMVK10%)k@2v)VynPx^PNHvu3x*VAj zC)VjSZ~lG>mwefWBYu&#Lb^p&Jo<%G9~Nnq{X1s*pr7S;)%`R zB@Vt8mkvST!yPPPg{oVqA=!&%g3)BQ8}To0!`m}gf$B48n5tK`gpEZnaQzW;8CqaF|F4PVQ>|Nc>V8~5iv3G!P zJ*LbH>_WZ`{X#N%jUxkC#F9J^=sDQ8@r_-!0(wQ@+2EtQ1Z7Oz2T1Z27~=sHvEXpm z`WUYBk0cx5GvSDh)CHFFuPwWKzjJ_Dy#PIM0rpFRvCmtp|9kEQ@LXyVZv~7+64=yM zpkoAy&qvTicOL&(SeT`sa1om(s3c1^4JawS*9BrnS(sn-Huwq=%OWt+Qi@t~0c|lb zDk?uIcO&SK7)uyiEnkHs8?4utuTX(Ljqvv8N!hT1S zI!aNXq(Ikg4I28g-u%HYeRr|l3dI^$SY(IPqAU#TbxAyWv`JVLO*lVQ!ME9o@A?w} zSi3Ln=ka;T8yTprr^k+{H80aTXuoiLT~1k;puWVFy^|0Dk0b!Nlo~EfjiW%{uKz_3 ztxOTC$LiOR|G3qvsub09{?}Yg+A_xazy|P$4lxux2!yHsd?Wi#GD;HD``G-en!|m) zmG!@KriV&B3vl_Up~J14DE{&W`4a#Z4FGPEAn&awx%WVni<8^u{^!Qe_N~w6c5kCp z!fB#Mk5GCYCN!qJbS+}71c?d%$CGM#l)5tRfQyLuqbUt!{O7|KFT_4y}#{OfqliYr9pMgW0QlsWKaz)m(r<9z89g{A|rJehl9W{YK=$5Dn3Ep_($+7tG^Rub3@_GHD8h@7JW3RczQV&Di4cdkia!!s%Y_PwnvKhzvpVO_1 z^xJ>-cbCuXv_oUyNSFE$6yTcor#QOfQ|kaoD8u0l!|ajx+U(5w7@8P8`aR-3T>Jb9 zc-R{*)}?p}b7%>8J6L)AYOgL_PuZ$_g9PKF7uVVR#l%ePcp`W~-!~wcyyds}O947i z4_sKU4cXb1keLo(Y9UtN%9~lHP724ij_7_Xy29WJ3ARP5oTT1HgMERzE6c#HvUR;S z>xN5rTHul8z!mNcy(cLC>O>1!y}>s+7(j2;&ovHtmLk={t--D!o+giybZ%<9aB+*! z=Q6&ZaMJ$O0muw6jMF;YJ$j^Ghc&3Q3DR=>wtG04Bhz$t&CL2>sX4o0(R3qU&r7SJ0FqWr*xDgm6R7;f>_F2^5*dhe?h)!)s>0rLb(jqDT$rL3uHIrDar zw0lyn1bMNJX(Wpjr!z~qC>-n3<((P^m3`4R%LMa`W!7QV*5U}8`-nca0ZE zqWSeEh;}`ho~Ev@jZ7lL(*a+z$FA;tD;w_KRW;|X`-(dEj9-fOjj??!V!ycBVaSWA zne+Zah1f;ejr&f}38Kb04;WgE^WIr7Bb-t@)i6Dtb6F)Ea{N=3cmL$n-y?hL(sk#? zx5tdzu7x;Nag>yG=$fHB+ZRqore$?GJe8dR8!r4AOu}gr8dy^~#@&9tT0;Z6(v*!( znyR_RsJ`A(x06_%(&RQd6t__3{8|6KU(rCZy4kZE)msbv)EVQIKlQ-xJA-PaoQNxN zW?drkDgy5GGY^vX(KZRZZWtT%S>_bu5W|>hbD(GQVoz8#rP$xy5j!A8-)Eog8$Obm z9@T#Z!|-{1%C5oQRsS;mxmLW@{@LNdh77wWhlt&-L_6m)+HbG-q2=jFyX4}!vVx7W zC*78mXC>nFt$g`D@_gM7WjMntx-jYJQr8({sh9t}0sgZ!2kzEGDPO}DcJ*WI3afJm zRmS+-g6CJ=56l@_<=!feWK)-se)C#Ao4FoL`!&VuPg@YU4eyiPTC*%~jKsLRZ{VKbm1LCFNOLIk}lj9fio zer!arSloEWI&fJWRhs=?=IxNwz%(c0UR^M0lW5E7je&#T=o?XiRi^GS!(7sK9h^7$ zc+$VQt6Jrf7xkM0nM5N^g*1JTpExu^dVljvpq8a6=}B7>fP?k&aykdD%b1(@c@fWv znmNYTA}h@{9y;>$Ft!+5U2MRh0wpqWa=c4unVGL<=VBJRC6w;p!<%_%h-;C_+4yS)=XDO_pMcj zU!3dl1J>N{c+M;7_bDFg=dfD_{jl>ps0n51hzKbMJI55AO^5iD{2itSm#o#Dd4%>( z8LVxWXYaqOeg?DEfuz_DKXv{#F*q)hJ-MT;;B_{^FbZVqzSMv;+(i2(-*OBgT#I>lkub z?ko=hxzT4&^uEW2>mji^01N|yIyeBm86+wCGPoFus2ZwY3kbls!0#lmN~#wfOVe ze@W8m>FNIYM%(`?NlKyMTlL9B?I?V-IBi0)i9n0T-GA^IIo15`+*rw(!tQgb?!PQ) zd(0?hLGF9PxjiAR2t@1ue@{gJ_m&X6&w1BTqLc;CdJ=~UU{!`^Z%42h5O^Us+p+8h zM#uftriy*=$Tr#HF5#%cLz|i+nZnsrKcMYu8FnYH?PZfts%l^B;cth2Kn#jxZ$#;!rQSQ0@Lh;g=O+$HRjv`iplaS{H1srqnK6)u6`Zr}F1lFD;fSo!3h>oV>&WHT30Tpn2y4 z!_7T2I5({7i@IB~a|n_OHfD8PuQ(2PGyP2BezN$rRigEgCYc0JA83%Is57o1huJgMkC2UW{0ID`IW-(D93 z2mVZj8wU<9_xVFy9uSZ?{Ty35O~qhwwiuePkywNrfnxVClt0{HkS=FFVr?LB>u2!W z+#RtjTXVCI%Q1q>$E)WGsOM`Y_x?1Wi-s8B#fx_)(o3ok#~f-hmb{A9Y6`p~x<3#+ zIzf23@&KpAy`Teie{xsk$J*lUU2eC%J|(D<4Zgo#8@uyTLdH68n(Cpw%g0u2J>E@q zrP}Hje&k2}t=rI}nK7}On3w1a_txo*Z9Q3uIq2V%Oca;soi4NFCm3@bjH9Zy|E7dq zRQQLkljd)7h!{#J$GffehHtZBG^W9%GMD;PahPv4Q&SWM1DNK;JwG_SMI}qXq+$1= zsV0`$TogN(xz4ygM@_Nzc^-eZb7}rJZ~EpvwmS)zQ;YcKuHpkt_+@ft+N|ml9)Vmm zABI;{V^3WIxUehO1v7mFN~X~xR~my^|N43{fo{tV&<7?I+PTg1SFB0-9N(tCQwkl& z^ftR=ZP(n(MZdzIn!Gxcq3DfJytl1|cFz|L<(_g90m@2-c9G77=;<&cN-zTi_1G@g ztmzdx__;>>nBi80M`GF{$AxLyZd~O0D-1-UD=2cG_VaO7wIQkOeF1DlusvukvSG8t z+li=N-08%3;QHcbv07*~Iqi=#X$$Zd5Mm-X#qg#0H|Tfb#lcAPw2Cs4#r@Ld+b1Nn z%W0Ack-Zu%i)SaAmFpJ~OkUaeNe|sD0aGYe-h0}uXLrNZ)v=B8BC;#nS-=&=be<`p za-UUntP`TC)g*dRgH=WQ7!RJ%dgrI_vqP3;{J3GR)6e`rSap1?nrAa`-03`E&2eIh z@o*#KOD-TL)_yajbu&c0u5! zTH?g4=sy1&g(OJ2l%2Ke6R{_m_S^7+#W^TE6IXJSlwNXPZm=(->-fQ_hlb8+B&cMW zD_$+b{!Db=*fmHeY8l^ze0w~GSW1Q>*Q1TgnU9jDyD&-n#z|4&yZrO2gMRjywWSiV zX?x~3x$*k~_8VP~u&O%)O5sKV^`FXNC?8Ip1x+1IEq(}hefpB~SJERnLY^-D|IE|T zJ_7P|k?r4rq<#s@wzNFv0my|(mYdQg=q)cPt|-83=7l&(tAER=X>z4MdY-<8t}(3% zeZLZs1`Rosv0UgbP4|I4#n`!(WERd{T9?wUVA9LfN}hj(&gq>6&Om?J55+>W<}cp@ zrn0`j4kgwC$2P2tl_@bGKIW`!l6VO40nU4b&erK(hlzZha9vzR|e zf3PfY^HT-9c2OTfX}j)c+XVG+tn*x^S(u>~Y-D#^v?;nlb_MvuUR4g``Xw#d*Z^tt|n1= z-rKO7Uzp~dW30nM zIQzJ}QT!6GVV_P86E?<|+B5gRpnD)H=a9H`da)>G=-^F38QreAkUgZf>ZpldI5&`e zhlEs=eS4&UEMjl!Dqe|lj<4;d?i)<1UtxzxXLht#%0u!b`;g;bnlt=%8_ue<)}a4M z%^Hi+8NZ<0#!^hidLK!I&BBjD$dh*WyH1m{%xZ36AI;DIF#9mo=trmM6f6wRYf9g4 z%Foaz@|)LifsPMpS;O+0(-DKre)2pI`Rh`{%?f82g)8!cM$C*SRUzk*OVA@=GXaw} znP=*Y_zkFlZ=6tft~@VD&u^NM1MkFZ%}Q;5rX@zIdkDq9Vngu(OOr3w?ozFYDVd+UIwhUcXqGmwR@)eTn!t}NtBP2pn+ZxE5yy1nX7|5MIBo$LrMa)#cSc)H)(t7u8*X9P^fW!m?O)}TvW zzlROY2X-UVp(iw$Zz=Y{beKbCn=I<6VP8e7%PEd=8QTrZ#QqA_mFM$vQWL}r>t>&& z!JKZVDTwn>DX%}rQP7U}X{7>Xtp!^`Eoxg`e%*^*3HvAUBhnO#6=hK=vBG`bwi_8f znpUS^wjpyt`Jco|O1KeM%lP?-k`)mLn@r11N))nR#WYAl##EM2PO`n+;r;n!YeJ!- zqGGK(oWj>Ai@rJ;W*#UI_i3Xa^QlY$`AL2KqRrjy^5njR_epDxy(F70wb;NmEY7VR zDz7m+Js1_uAOUZj%&AHnIOLOB-->W{vl_9H=zJUcz^q-LDsx?0BslXyPNO2#sL5l} zAQP52j#zPx96tX(&GbOavVj#BczT2Yr%KNzQf0ZTPX(uJ({;Y;I4}H5bzXQ+XC7Cp zX*C9B9;!jerkK_Sv50bfH0XM6QepOhM{faf*d|`09cqC@6kn8o)3)gL$&}ex`+1dH z+@fI?2UN0m@pDHr6zFpW2k6Dk>ygy&2wXoMQn5>qrhS|=cev9 zc{huD_nVEb(I+~gMx!{L&CI|z;SWr^)HPBJ=1+ZpUSP&OHC@-2?AzSz(oHJCyyAm4 zgCl&%{%X-j#l8c)M|~GTZeRs-yJxZK_zafwzE1cNHd+2C8&jq6`^Egtwl>e4i+ZQJ z3)@pEH~$|rZLfvYYee+=X9B@%3wfz-Fo%!}ZAo5qML<26tTZJJrzvb2Ll)FBL6BXGOLzCFqX$9V$4@%>Sp3Vb|sTBu^uybwWStgCyQ>uklp zma?Gh;)o)g)~N09c+73TTdB=&@Os`O&b4153U2nOD{Za*wK=tO?3=GD)=y@&C13@H z@dCz=M8{x!Wp``R8i&LNYwku4liFB8$tLW6^q?%Kw54re97q|Td+Z&zbUDP*e zLB{-huPf`6+wZ2zeCVRK?8A$BPsj zk}$k9y!epmLRY-`oE4urvv_?R3$wob$C7~R$z z=eGMUQ{sh8I=(8vw0KYI7|(hH#ti>CP3HVhQAU%f2>K(3dF2;pTvIKq<2S5NxNR~H z6|a2xtiz#ly6XrYNYsu#Q~a2nuBed`y_dIZ8ZoOn>35t!9{ct~5iFKTzW@`XpCQbA zI}?f6Eo(beNQ&lWBb|*lad0zFDTB0pG?Ds!Dl4u3>r{DW=x+%!^Ayw8e|`q@bhujb zgkjZFIs;eNAB%;ZI)*waN5mPXUlVrT|Mj#l0C7>)iJJ8_#6)D4-C_gF-r~mSTOmA> zV5*8#F1I9L;I%@-vqwoR_;CZyv2o1)ZnDt| z+G)Wqe%CFD0^5meY@^E~s;qGEVtHfW}OPv_jNzT82D2(tkrS@afX_-^=k8t+< z#I&Jk|0%OKtH35WU?iK`l$m0fd_1Rb819UoxzEA1^Bg~SfA?|y-=0E=`f*ggpcE;JMyih)$jT-JEHz2s!~WC=1HGKGk0gO!*lvP_T7?<&R`+W*oxhSgU>{n-4^T0ndTHk-f@XJGXUSd`%%1~;! z`2@DQFTBuA@3=k{*atoEeCN|Rwa0xlB_|wgpuQ^hQ@6_;{6VTy^jSebEna*IQTK&Y zuRRCdD?;~cQ)WHf!TEIBv#|OzgPZ7HTHIN8ZN!gjyD+V)pN0x$Ms*);PqNU=Qov3xM+>rp z&P$K&ikGBoK#pVOJ!k^WO7!E6Gt*XDD(MZ6S2}I?r0&TOPaWoYTqQngUtT z7fiL`l|l91*|J~w;Yr5V$qMGJPHB>02dte7 z0vf0%loU6k*k==7Va4q;#u;Usehi{zCskj;U_(WIMk)f;J0(EXOv|p4; zXzfZtd0@*4WIflilhog_3@_~tqp!tP0!AJs(%L3Ia2oZ2j-P>ba^?tQ1<$y+B*$ z9plg!kc*IeHcNGGlPSWoU29?Ybu#{Z;jSQk{OQ4{+_Q(fxQf`7Kc5qdC*D=@d7Zn* zFvM%>`=sZIK5*=6$WAZ9H)r-R!Y~N)?$vsTZ=PhOcB zmqh*35c_PQLw`Ar%uh$*$%DlX17`7`qsioh@fWYA9fP)#=rzRIacNFe1~e zo1c-hVJ=LCGXH$<1+24r3^BF6WvP3D1#R*c*x4aY%cmxvRwhVJ@-v=qJvv|^C0ct1 z;%bkV`EDp6zH6y`(ohFA+bOYyG3qm`7K%@7nU*bR8y<$fY$NxPQuYtg1S_3V=}~@I zFs?{=RzF4*FtF?J85lKiel{B^t^6R~APW;U@pXr}42kd~ZOb&s4TeqEa9(Y{SU}w+ zdjA}#T&LaEAwP*Xa(3=?o9hr%&iDV4K6*MIxzJuNe#PE^DDZgzO|8d7G z|Fh2?Oc0U(>e+!pSK}LfFRK}IZv3Zs+|%`6W=vruIJdAs`+N=9gTlsov}+8d1mCEx`G@*BgMbX<5G63_SO zMeIukje%;RfWPm+-Mz2f$^tO!=}vvmx!ajI)qzgF_+)T+S1snX2jghIERwzOV)|5^ z0|T;sJM<-DyzilZRrB+|tD$0HRQI0BnfYzAzoJSfedr_i&~K|?+%;v18;3nh#~ySD zcsyWl8N^?V;rqXaTvOD)oV~@AnX8zJ6y0 z(0n{bzU2|Ci0q#;QCW}3{S5bi&7y4y|8-NTU3z7zVEDl;J@PJkJ&`A3om#!q9U=XY z^H2)zpwl79*v*Fu5j91NTz_~WaOHkXa}ek_yArA5unQ|4&hG`>JlZ1PjrALlQ3vMI!)SW}#{N&;AWA)Lo&GtDVK5_&xP_}bH2a5qd>)jHWS zChDWBl*vY}`aLaiVsEtFzDOGOgGl)XKT$u-6`50QIhT9~jd!|jR_`m%eyIJ%@W;nJ zii}#l)0G+~sfGS=rx(cQf1V`NU7`{R&TSl7H;gStF~;|8kAgjX@xqQi!plJ!!qJk`wYtJRp8U49(vbcDS5o0Be$2e!>Xk~*dw>)i{odHT6m!7hhndCubkuWx}OeSh(SFdwZk z2$$afn^i>c%+#dvpJRT7i%O|K90S)uxOZu7q8>X)o z%)6c}m#N2pk4|%BL|wq0R}4M;Tyc3|G_BOJ^9-V z7bv$D`gvT51mCdOdSXoFsR<|L^?<><^0u$5pLzrW1J1U)LnY(5fcOrvBusI+xrH=jut4 z(7W$z6v&h`z(5}hP`R)4qI!;fkvb=JnUjB@qMc62ZdAkfN7$r}flI5mjYHH)$L7Pf zvdx$F(UkY!m?5XeVpS{wt+tmz)Q0YOprlL>M^vHDfb%R}X|b0KE9Y|;{v8)4rUJGG zPRY{5Jtd_`MmNtfnX}Ph;jyO*%WU#^R#NJ(QStOvx zy&HWgSM1MHY6>je59ru{SkRVgQBeaWwCj=C|7%kz4fe$CcaFcM)RP_C7NtGp+0vBG zi-S`$uV;AgEB%Z>fp^tD>1aRQCvlaT2@1%|@snIqr+d!KcR{9e^$Jd__CQn2uBChA=H0oC5GwhjJ^P6fpG0aT0PeA)3^ zgR1qIaF#!4LoRq3p*=AVTZj{7>`4uTG? zU)?EQD1iQvmJqb8_rFtf*i%3G_fq^o&qkX>x<#h*$aWjoRUe%FcGlvL`{n*!^T=`} z9ks@a|8Tp5^UMj<_f2kOPOQ)Go@<63Py;7sp+#P?8ifxCO^!DYMdU?t^;|xtXu1h3 z&`uuhg6uSr?_8EUJeQ|p8&SD?fnw0%kU;4wrM7DQja4+dw)U$e1ZNVx-G)Jlv%q29 z@v=lrF+fYP$Lf-vuhZmfqU;hBf%9n;l z)_@ud9^3mo`QC~mr@Rl3p1f=F;SLmaT)t86DS7Fe5Z58hCu{!ch?Hb-9~tbe(2aBZ z+o~(Azh!34LBsI&NBDJt;qY21f#*AZZbQ4XnRwwSR8*kYwlw)7A4aPC?}tR$$IagU zZNYt|KSwm3GEw6g$s{A#a*RlPdem*$mYEfs1>SRL2ZnYbZGjjgOy4mQ&%U1Qviv6V zvm@YwljzcUoB04Jt@2szSoYT*m3==7^oGp(2g=)P`=a}t2?fe$@zX&#n9{@iCT8x6{ebs;yX$Z&#{WB5!5% zLHC2UbBncS&OFPvaLNA4ekH-dVjc{}-Nh7s->j)4lK*ag=)oBy=_vd{ei0|Fo|T=) zPj~vBl+H8wEHuJsp_=Dok%p37ZNvV?R_l%jl=e&Q4&4sBX7g*d~p>mmRRH$lufj7RS_|g@NXTQ=8NWYCO#RwWD9UHfPMU z*UP!H70L?OGi$pkboI^deI$-69#5oU>3lWZ(Zon5v`>CVzjJaWhb^Tzr)wZMo;*yWCS@xn=jJ-)+cMNfa;k`dRf_HCT z+_8K6)x8#X2fT2G^+Ddp=G4+s>&_j%&veTRK-EtM*!t_rTr^yGiCx&#W#m`x6rXS3 zG+qsu)Vug5N5Ayqa4ekofs1shUjkk#6q&yxl|1iS@9Vxh7gPgVj% zkX|z@c70Rbe}rCAOQX3NxNH0ak?w-`yclAZg6|2Dmj28S>h{K5pG|VXg)=EgD?3v; zVZZ(#8aB#+)(KZzS=llR_YXSv!jzXvkKI(S`eVYLr8Qyy<9^3{H(Z>1mHj$fU*&t# zL%m`7gU0$Z%!ij%lTdcCecj{R*&3vgVh(-u%E{;&>aeyp<=dmbOoyirk@T+LP$w?OUe2{)`*bE79d*Ab-8Y`j zz+PH>n_@fNb`hgE+AnGkZhjQj`=KRmPTeqMf35ci!m`_&`0)c?_+g0sj|ePW3VN{b zO$Q3$Cp7I$+P;^OwTv9*SFPdU#~o#NT7q^ymT*5*4ebDmlv>utx0shITV;b_?AR6D zD+fF#kMRS4S#<%$c~+d*5apkVp+CIlXZ(xTMEoIG(jfbK9@YNky)%(*45?edpp1u?gG`h=_=qephL1~?Q>4+99|B}C7 fL#V>b=*Lq>zwLUU&2=;gTob7%X(-l!-hB8k(o+$u literal 0 HcmV?d00001 diff --git a/src/main.cpp b/src/main.cpp index 90fcea1..600c29f 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 24; // feel free to change the size of array +const int SIZE = 1 << 20; // 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 d2c1fed..c4ca0de 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -13,6 +13,8 @@ #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) +#define blockSize 128 + /** * Check for CUDA errors; print and exit if there was a problem. */ diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 25b63e7..c209dca 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -3,8 +3,6 @@ #include "common.h" #include "efficient.h" -#define blockSize 128 - namespace StreamCompaction { namespace Efficient { using StreamCompaction::Common::PerformanceTimer; diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 1d12c05..4cb0015 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -3,8 +3,6 @@ #include "common.h" #include "naive.h" -#define blockSize 128 - namespace StreamCompaction { namespace Naive { using StreamCompaction::Common::PerformanceTimer;