From 511ac01b8dc0bf8cf3a2f3070710b544ccf2d206 Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Tue, 17 Sep 2024 11:58:36 -0700 Subject: [PATCH 1/8] add cpu implementation --- stream_compaction/cpu.cu | 43 +++++++++++++++++++++++++++++++++++----- 1 file changed, 38 insertions(+), 5 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..b766037 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -19,7 +19,10 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + // assert idata[0] == 0 + for (int i = 1; i < n; i++) { + odata[i] = odata[i - 1] + idata[i - 1]; + } timer().endCpuTimer(); } @@ -30,9 +33,15 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int count = 0; + for (int i = 0; i < n; i++) { + if (idata[i] != 0) { + odata[count] = idata[i]; + count += 1; + } + } timer().endCpuTimer(); - return -1; + return count; } /** @@ -42,9 +51,33 @@ namespace StreamCompaction { */ int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int* binary = new int[n]; + // construct 0 / 1 array + for (int i = 0; i < n; i++) { + if (idata[i] != 0) { + binary[i] = 1; + } + else { + binary[i] = 0; + } + } + // scan + int* scanArray = new int[n]; + scanArray[0] = 0; + for (int i = 1; i < n; i++) { + scanArray[i] = scanArray[i - 1] + binary[i - 1]; + } + // scatter + int count = 0; + for (int i = 0; i < n; i++) { + if (binary[i] == 1) { + odata[scanArray[i]] = idata[i]; + count += 1; + } + } + timer().endCpuTimer(); - return -1; + return count; } } } From 4c7b6ecd3c02c8270203452e1b4acc6961fce838 Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Tue, 17 Sep 2024 16:59:37 -0700 Subject: [PATCH 2/8] add naive --- stream_compaction/naive.cu | 56 ++++++++++++++++++++++++++++++++++++-- 1 file changed, 54 insertions(+), 2 deletions(-) diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..b325ba9 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -12,14 +12,66 @@ namespace StreamCompaction { return timer; } // TODO: __global__ + __global__ void naiveScanKernel(int n, int offset, int *odata, const int *idata){ + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n){ + return; + } + + if (index >= offset){ + odata[index] = idata[index - offset] + idata[index]; + } + else{ + odata[index] = idata[index]; + } + } + + __global__ void naiveScanFirstRound(int n, int *odata, const int *idata){ + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n){ + return; + } + if (index == 0){ + odata[index] = 0; + } + else if (index == 1){ + odata[index] = idata[index - 1]; + } + else{ + odata[index] = idata[index - 1] + idata[index - 2]; + } + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ - void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); + void scan(int n, int *odata, const int *idata){ + int blockSize = 256; + dim3 fullBlocksPerGrid((blockSize + n - 1) / blockSize); + // TODO + int d_round = ilog2ceil(n); + int *dstFirst; + int *dstSecond; + cudaMalloc((void **)&dstFirst, n * sizeof(int)); + cudaMalloc((void **)&dstSecond, n * sizeof(int)); + + cudaMemcpy(dstFirst, idata, n * sizeof(int), cudaMemcpyHostToDevice); + + timer().startGpuTimer(); + naiveScanFirstRound<<>>(n, dstSecond, dstFirst); + std::swap(dstFirst, dstSecond); + + for (int d = 1; d < d_round; d++){ + int d_offset = 1 << d; // 2, 4, 8 + naiveScanKernel<<>>(n, d_offset, dstSecond, dstFirst); + std::swap(dstFirst, dstSecond); + } timer().endGpuTimer(); + // setFirstAsZero<<<1, 1>>>(dstFirst); + cudaMemcpy(odata, dstFirst, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaFree(dstFirst); + cudaFree(dstSecond); } } } From 25f31cddf700ab12be8c917d4298fead1f468495 Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Wed, 18 Sep 2024 17:48:09 -0700 Subject: [PATCH 3/8] add efficient scan --- stream_compaction/efficient.cu | 46 +++++++++++++++++++++++++++++++++- 1 file changed, 45 insertions(+), 1 deletion(-) diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..921160a 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,13 +12,57 @@ namespace StreamCompaction { return timer; } + __global__ void upSweepKernel(int n, int d, int *data){ + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n){ + return; + } + if ((index + 1) % (1 << (d + 1)) == 0){ + data[index] += data[index - (1 << d)]; + } + } + + __global__ void downSweepKernel(int n, int d, int *data){ + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n){ + return; + } + if ((index + 1) % (1 << (d + 1)) == 0){ + int root = data[index]; + int left_index = index - (1 << d); + data[index] += data[left_index]; + data[left_index] = root; + } + + } + /** * 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 d_round = ilog2ceil(n); + int full_size = 1 << d_round; + int block_size = 256; + dim3 fullBlocksPerGrid((block_size + full_size - 1) / block_size); + + int *d_data; + cudaMalloc((void **)&d_data, full_size * sizeof(int)); + cudaMemset(d_data, 0, full_size * sizeof(int)); + cudaMemcpy(d_data, idata, n * sizeof(int), cudaMemcpyHostToDevice); + // Up-sweep + for (int d = 0; d < d_round; d++){ + upSweepKernel<<>>(full_size, d, d_data); + } + // Down-sweep + cudaMemset(d_data + full_size - 1, 0, sizeof(int)); + for (int d = d_round - 1; d >= 0; d--){ + downSweepKernel<<>>(full_size, d, d_data); + } timer().endGpuTimer(); + // Copy result to odata + cudaMemcpy(odata, d_data, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaFree(d_data); } /** From 884406fd69955201fdff70badb9ddedbbe39dbe2 Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Wed, 18 Sep 2024 17:58:39 -0700 Subject: [PATCH 4/8] fix cpu implementation --- stream_compaction/cpu.cu | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index b766037..acca3bc 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -18,12 +18,12 @@ namespace StreamCompaction { * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); + // timer().startCpuTimer(); // assert idata[0] == 0 for (int i = 1; i < n; i++) { odata[i] = odata[i - 1] + idata[i - 1]; } - timer().endCpuTimer(); + // timer().endCpuTimer(); } /** @@ -64,9 +64,7 @@ namespace StreamCompaction { // scan int* scanArray = new int[n]; scanArray[0] = 0; - for (int i = 1; i < n; i++) { - scanArray[i] = scanArray[i - 1] + binary[i - 1]; - } + scan(n, scanArray, binary); // scatter int count = 0; for (int i = 0; i < n; i++) { From ae01c22ac9f36043af93a34f1beba02230251ec8 Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Wed, 18 Sep 2024 18:24:30 -0700 Subject: [PATCH 5/8] add stream_compaction/efficient --- stream_compaction/common.cu | 14 ++++++++++-- stream_compaction/efficient.cu | 40 +++++++++++++++++++++++++++++++--- 2 files changed, 49 insertions(+), 5 deletions(-) diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..0acc872 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -23,7 +23,11 @@ namespace StreamCompaction { * which map to 0 will be removed, and elements which map to 1 will be kept. */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { - // TODO + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + bools[index] = idata[index] == 0 ? 0 : 1; } /** @@ -32,7 +36,13 @@ namespace StreamCompaction { */ __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { - // TODO + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + if (bools[index] == 1) { + odata[indices[index]] = idata[index]; + } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 921160a..1764c83 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -40,7 +40,6 @@ 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(); int d_round = ilog2ceil(n); int full_size = 1 << d_round; int block_size = 256; @@ -51,6 +50,7 @@ namespace StreamCompaction { cudaMemset(d_data, 0, full_size * sizeof(int)); cudaMemcpy(d_data, idata, n * sizeof(int), cudaMemcpyHostToDevice); // Up-sweep + timer().startGpuTimer(); for (int d = 0; d < d_round; d++){ upSweepKernel<<>>(full_size, d, d_data); } @@ -75,10 +75,44 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { + int *scatter_result = new int[n]; + int *d_idata, *d_bools, *d_odata; + cudaMalloc((void **)&d_idata, n * sizeof(int)); + cudaMalloc((void **)&d_bools, n * sizeof(int)); + cudaMalloc((void **)&d_odata, n * sizeof(int)); + cudaMemcpy(d_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + int block_size = 256; + dim3 fullBlocksPerGrid((block_size + n - 1) / block_size); + // efficient scan + int d_round = ilog2ceil(n); + int full_size = 1 << d_round; + dim3 scanBlocksPerGrid((block_size + full_size - 1) / block_size); + int *d_scan_buffer; + cudaMalloc((void **)&d_scan_buffer, full_size * sizeof(int)); + cudaMemset(d_scan_buffer, 0, full_size * sizeof(int)); + timer().startGpuTimer(); - // TODO + Common::kernMapToBoolean<<>>(n, d_bools, d_idata); + cudaMemcpy(d_scan_buffer, d_bools, n * sizeof(int), cudaMemcpyDeviceToDevice); + for (int d = 0; d < d_round; d++){ + upSweepKernel<<>>(full_size, d, d_scan_buffer); + } + cudaMemset(d_scan_buffer + full_size - 1, 0, sizeof(int)); + for (int d = d_round - 1; d >= 0; d--){ + downSweepKernel<<>>(full_size, d, d_scan_buffer); + } + // scatter + Common::kernScatter<<>>(n, d_odata, d_idata, d_bools, d_scan_buffer); timer().endGpuTimer(); - return -1; + // copy result + cudaMemcpy(odata, d_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(scatter_result, d_scan_buffer, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaFree(d_idata); + cudaFree(d_bools); + cudaFree(d_odata); + cudaFree(d_scan_buffer); + + return scatter_result[n - 1] + (idata[n - 1] != 0); } } } From 9914b079e01a1b420d1d0eab159f1e4dc878aee4 Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Wed, 18 Sep 2024 18:43:10 -0700 Subject: [PATCH 6/8] add thrust implementation --- 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..4a7b222 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) { + // Create device vectors from input and output data + thrust::device_vector d_in(idata, idata + n); + thrust::device_vector d_out(n); + 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()); + // Perform exclusive scan using Thrust + thrust::exclusive_scan(d_in.begin(), d_in.end(), d_out.begin()); timer().endGpuTimer(); + + // Copy result back to output array + thrust::copy(d_out.begin(), d_out.end(), odata); } } } From d09a16df6a5374e0f4b736bce72f9531efbb57be Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Wed, 18 Sep 2024 18:58:45 -0700 Subject: [PATCH 7/8] update cpu implementaiton --- stream_compaction/cpu.cu | 8 +++++--- 1 file changed, 5 insertions(+), 3 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index acca3bc..b766037 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -18,12 +18,12 @@ namespace StreamCompaction { * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - // timer().startCpuTimer(); + timer().startCpuTimer(); // assert idata[0] == 0 for (int i = 1; i < n; i++) { odata[i] = odata[i - 1] + idata[i - 1]; } - // timer().endCpuTimer(); + timer().endCpuTimer(); } /** @@ -64,7 +64,9 @@ namespace StreamCompaction { // scan int* scanArray = new int[n]; scanArray[0] = 0; - scan(n, scanArray, binary); + for (int i = 1; i < n; i++) { + scanArray[i] = scanArray[i - 1] + binary[i - 1]; + } // scatter int count = 0; for (int i = 0; i < n; i++) { From 246463ce5b250d10671f661f56b55b59e0476012 Mon Sep 17 00:00:00 2001 From: hazard-10 Date: Wed, 18 Sep 2024 20:01:48 -0700 Subject: [PATCH 8/8] add readme --- README.md | 62 +++++++++++++++++++++++++++++---- img/nsight.png | Bin 0 -> 140196 bytes img/perf1.png | Bin 0 -> 150048 bytes img/perf2.png | Bin 0 -> 147521 bytes img/perf3.png | Bin 0 -> 148138 bytes img/perf4.png | Bin 0 -> 148955 bytes img/perf5.png | Bin 0 -> 19682 bytes stream_compaction/efficient.cu | 4 +-- stream_compaction/naive.cu | 8 ++--- 9 files changed, 62 insertions(+), 12 deletions(-) create mode 100644 img/nsight.png create mode 100644 img/perf1.png create mode 100644 img/perf2.png create mode 100644 img/perf3.png create mode 100644 img/perf4.png create mode 100644 img/perf5.png diff --git a/README.md b/README.md index 0e38ddb..2c56049 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,62 @@ 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) +* Jiahang Mao + * [LinkedIn](https://www.linkedin.com/in/jay-jiahang-m-b05608192/) +* Tested on: Windows 11, i5-13600kf @ 5.0GHz 64GB, RTX 4090 24GB, Personal Computer -### (TODO: Your README) +### Project Features +* CPU implementations: + * Simple sequential scan + * Stream compaction without scan + * Stream compaction with scan +* GPU implementations: + * Naive parallel scan + * Work-efficient parallel scan + * Thrust library-based scan + * Work-efficient stream compaction + * Performance timing for both CPU and GPU implementations +* Support for both power-of-two and non-power-of-two input sizes + + +### Questions + +* Roughly optimize the block sizes of each of your implementations for minimal + run time on your GPU. + + With array size set to 256. I have tried block size to 256, 512 ,1024. Of which 512 block size showed best results on all cuda-enabled functions. Showing anywhere from 10% ( Naive po2) to nearly 250% (work efficient po2) improvement over blocksize == 256. 1024 Block size showed results on par with 512. + +* Compare all of these GPU Scan implementations (Naive, Work-Efficient, and + Thrust) to the serial CPU version of Scan. Plot a graph of the comparison + (with array size on the independent axis). + ![Performance Graph](img/perf5.png) + +* Write a brief explanation of the phenomena you see here. + * CUDA Work Eficient being not efficient + The most obvious surprising result is that the supposedly work efficient is significantly slower than Naive or thrust. Diving into the profiler it seems to be due to too many kernel invokations. Both upsweep and downsweep have similar computation and memory throughput to naive kernels, but consistig twice the kernel invokations greatly hinder the performance. The gap could be narrowed with much larger array size. + + * Thrust implementation + ![Performance Graph](img/nsight.png) + The memory throughput is between single naive kernel and work efficient up/down sweep kernel. It invoked the block size that match the input ( 256 in this case). The primary performance jump should come from a single kernel call and reduced back and forth communication with main thread. + + +* Paste the output of the test program into a triple-backtick block in your + README. + + Config: Block size 256, Test Array size 2048 + + ![Performance Graph](img/perf1.png) + + Config: Block size 256, Test Array size 256 + + ![Performance Graph](img/perf2.png) + + Config: Block size 512, Test Array size 256 + + ![Performance Graph](img/perf3.png) + + Config: Block size 1024, Test Array size 256 + + ![Performance Graph](img/perf4.png) -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) diff --git a/img/nsight.png b/img/nsight.png new file mode 100644 index 0000000000000000000000000000000000000000..c3375109d759f9a5194778925fda3dcaee96f256 GIT binary patch literal 140196 zcmc$`WprCjwyrC495ctv7()y(Gc!Aen3)-3kTGV8V~$y7W~P{#nG-Xk%yg7*f8G0b z?>_gAJI;@@#?X*zX-QS8nzgFteBP;WWko5Jw*+qi004@Nw74n&fXEI2Kr_C9hn%6h zsvLwoLAj_(i2^Fdi4GtISPKya5dfe%7U{_d4pK&Rl-70u0C1`QJfLJ$sm=j_*L)dq z5%n+n$L)HS#F|OB?bB3T)oG@@;d39fNT+(GNu-qEksU-a$Er;NMYQlTY2NmR;dZg! z1PgaK5_CpH6UEVVR0-p;hO-jzbBgDz;9NeDA0K5~2TOD$o?gxj!9hFpXe*vA%FK$*1Qu0=s)lJtKOS^(9HarBhan zv}}^g%L;qNy5!&XJ? zdZPpF+gt9rwkrR0V5^dAM)->TomYvA>j-;6bd{X}mX%4)b+?;h1zwFr`$qj2!lbVi zCh{4mdhG4b-%$E>SJJW1(*}7lt{jRkcN^;Kr^-J!o+AUSf6UWzPW6-ezR%L!+RxI>yO2*RfWMjQ z{jQ&qaEkFGg9OxYjEzD>$vc+Sw|PHB!n?vFHzQ_ud37G8BOouNz#O9%;@9aA2y)+8 zUsv1!-ezP5p+FnlmgP2dA|>htP=xwenhk_mMs2?!q$&^ZR=0e(B~RU$2X>kGTD(r+yg&DJ>-I2 zL>#3caSiNfSC~4SM0Y~^h2`*C)ult;oK=OP(xJJ5w_^JQ+~s<}vi*PuMcOWwzQG@E zARzs`v`M+pe^8j^3koIm)nW;Ic7QPUkCMIE1Ql67#38?surzxwdbks(xq3( z$pbj5g>&?=2KI8cdS$pj*DN?jzz1#IA@j)XheYm@Z&UBoSlaQO6n)w%&6|GB(>2-D zDyW=#$uQfuB6QTy#H@hMaGYRuk;mscF7b-p)LN}0xX4YRClr#v2455$7*Cpv%u4gAJi?hJ<9Tiv^fwi}zT_TGB@ zS}d#R<+X915^z@Tkk1stJ3EE~sQZ_!_$GaFIEV#s&@aipgL*IkV5+$j>i)5N2F>#{`}+-R-r&Ow-O7C|^C31@R8r;0baS5DFvl}@hjLvOk*bl?rg*+qS_VZ5cc z8+v)XjPco1iwHfHMF9)HR~cUPvSC!iD0ksEBGLJZwr~s{KpFy0d6~{r)a^&Q$X&qI z?Lo53v}Gwo^Dn%@c+1zF(!?Nm|CmbAkyf9P^Kx z>|k@xw4a=v4?k6KhHC4hhU+*R9lBz_SIys!G#Eyenqj>)TUW0bACZbmJG}U@yyPCn^WE#`DfH05~`sKjk;&DBuUUgPgI$ z><<~|-%P$je@mBLG$iUKiB%}}b?`|4QS8k^&ouKfFbUEBuq$rTt1ZSMBqTbUzBZK1 z0vfM2_>3vzvODHYHNVUN%W8~M1d=NlWg;9}c&1v> zL2iyFlyOyF>P0R;wY9d61oJJ>S;NJ`cyPcfT*rX?O- zo+{6+Z_;5_l)%O%4Ws{d6FLJeQASyyJ9WO*ks zEOO*7*P0b{ho4rM6O?py{En*y`$;x|Mo|Pug{0v_u$p)=htm9Qnc{m&v&e{e-1Hkc zobMVTn-BYyBGJhz6F>o#^$Y|742Xp*AB;mXjTn}a@d<$T>0(n;THGfI zPI(V32~YqWT8Wj3W2aT|3D>YKbZmTWc>fm4ZAM^7!l`BAr-+Zx;7$NlAtt95?3%a7 z$P0`23B!XBpk&bcBMLD9K=8F!@R64mjx|1p07Ev?I6YH^p0>sEZcP6mT_S{M^0;dF zvOj$900lqMDp654HnO2NF79SB`k8@Y1}_=!pyX-0CN+lWcY>?H#oI|@0ba9=;otZh zbcc3M1pMP;^ud}BrK&OnsniT}G340VVuk!iz~B%woeWRaE*AA4L@}!5xw7JU62qb4 zU$Ba$SP$abI7E+IccqE&z1&iW`;qco(+=;xvbOtOQ?{gpN;>WyA)HB+d4T}9UD zdPsrppd)!~-na5E*ap;?mXuH4v0Ke=={3{TV4GnL8HN>_zcS6y3|x;%Wb&$i{7!=K zkPX0zrGXdjnNo_6jI86%REZg-GKR+$3fmZs5szlh%HiQj|M)Jm(|envx{`xqQ`bT6 zqyr&F9Z(|a)zImB{Ua(&BsFN;L6!)OG}6!)^#ow)m{c+S(G1pgvZnB5$SL(u;hjjG z0O%Din_s;bzSYX~6S#EqamZh^Bn$OD`?s>BE*spqA4@#6j+Sqb+~o(V=q?r$g3whSfxcq^N7nadU?loEZuH$ zd>DLDgt3eX_*rEVse6ob&{n~A1wp^VDJkijSR+wb7N$Ydb5pp>^6jdf~`;wx7Ovr5j$T*7uAk(B-p)kF5> zwFrQNkkNcnk&uR+nfSvChYcc%&%WI1l4L204^@w%;C%yX>19;t z`)hq2mY}wE46k4aPB2lEBy)?Qp{SSV?bBU_+>m=6>qJbv3tq2Irdm2Pc;wg~bF0Rk zCO(ZhY_0&lQLw#_ytua_%{xKzDVM0&?Zz;by~QYqKJn8K#RMsU`^jq=b?$PZ)!)S_ zqx`-LxwY5y{Z+okRL7LH@s4Wi$SqNE{37|i&?~&&LoSOb8|S8DZD%rs&Ar&HYw|bw$wV@q!E}#)6*H1rLH1?gg9?PS1G>)3K~UnJWMuL zwEdOGhg@L>w!RjgE7OX~nZPH71FL-o1@)Dwk542{Ge%hL>t;LHf(LnAS0GDnQAFZk zN>#&DSQfY4K4ox52fGQC-tyiF2fy#^cmJ*Bh|AwnSIUD|*q;e%`!GRF?fjg(ELXi$ zcrX~e597Ld)z!b)s#Q&w;A7c;>|bts?hv8nWgq3T$Xoc@I{R zmd>wiqzm02l&b%U!#hmB*6Xx^tP9(@>R@m;h=<))A@7AZlSb>q#@pLw5HgnZ`@OnDY_i zN(PV#0N`Y20E_uBMCBtcFR;s6Gl@Z57Oyt}q<|uWVXH^>9s!Q*Hfcd3a_cPWCO~#v zDtpz3KPWm%oT{$;aNsOEL7gxS-$!)9=h3_rF?CTu$q=T(G2vt& zP1>p*pooieW=8d0t;mwMvxwRSSWPzNdplmT87(!@#`#?5YGULWWBCl1+csB)-Tc9c znjc6mX+z&Bd7a6^JhyGO5k>GXUs=tsg3T!O*9j7WrJwXrUV3TVM&J3}R_sV;tia8? zXUMC%h6YKY*E#%VWPq~w)8&K(eVT?xS--E}-EV|@E?7QUmYSa`U3NY9H<29kz5%Gm zjl{L~gf3~!M2n>KA-{^qO#b+`g(s%w7Kai9c_}B%>3JDYDgip{B-W+fcEr_W;cj+M8wYj>?skcB zFVTyJRjZeb1}+mSw>W}bm&2Z`I8+Y67`m+e7VFBw356_|*=uw4_9=GxqH*5GU0$d3 zeSwY;^R=cU^hSr~Hi4RSu>mX-1QoOIbtL^u5xVyE80%gL^ZN>d2K74I?UIx3cgusK zm;-yK1ND7HSGDu3oGG=8<-~ua z>e}qqp1>OsbJE?$f|z|aN?^MvCpyai-YGsxyt#3IIBh)x6WO8z2pcY5pU=N#+AtbE z&gI4(VbGB*8vg#7B+L_lQXogkFBLHDk1|il1n!91?1hf$+nlzZW1`6&gWPPYerzTV zU)`o|>?>7uEf?O{p2yj2foGJ2rh3qX%94BIK`F(~RZ}0(*`bt4v8l|VamdE4P-Wq+ zn>mM=#sPp?_Ag)HuVH)TA=^@P_5f!96nq#orxte=wX&jN&;c|n)}NYnq^qBB+V5mw ziajxN`sC|m#p0IS5AniM55t#hyF`2Ci}O-pAW^1&JA{L#YqSe4+FTsO#<+0OilS${ z*Zd}xMQM~RW}8t&M`8c%dPvd#@?w@|nc>CsOhEZ+Ok54AS$Ir&kMj*H*8jW~Fg&mj z3EnK`@p`lo*pM%)oIKiGHRxRbD$1Ce@PPVpEH2vX!iSyKvyPc8!yrHtM_}{-)Ji2u zv~H~M#ZP!7ga+7QMB5K~h7$nJH^MD{hX>HH$Nx1VPzdR=en$YHJ{p6Kwc#0B`7B!y z1|aNKO_~V>nCog!4n={UB>H&+C5^X81~`BwgeB91RF{Wji=NPZyWxk%iq#VW11MrK z?97f62EqgU4~K{IdXwy7QU*hJ3t`78JbWYZ;{>qzgnQo(0SL#f+=NgO~5)c;pZml z3lFgirPXvWFSkkO>IKw4hRTUF^8y(_K>R4ki*niiGY|z@nYUcY78EV?#HsPwBj7|c zRHyAGo6{c!7xfr10c+8){`biJ>+6@ZoUn{r%L4d)IU#OLktBr)iKcLml^8}#7;`7?? z4;KhH*wRb8PUcKI7>xVn#ialOX4k6!rov+rju*4t+9jL# z%AN4jgyNYli0nFpTj0>`^6$7m5idv@OiDZVlOBA@>eBNvpqje4c>RK5Edej z%~&d-U4xHv!>QeX1MJ#A2!5Ga6TIV=V}H;L_bdon?#*h5v6$tji-tc60$9S|beL~Q zU8x@e+uUMz6)vF9ezr6f!em>kci*gex9K(letWwi!$K)Y?sqg3O_ZLxd`j1o;E(;C zNz6gG)6i99VZ^R|gLC@YjWvqlZ!F`>a%i-|oGE4V!=LnAn$~V+BrPyW@nO;1-p2Wr zc2Mq4dqZBujdnlM&o5M;h@yF-1EUC$Rqb|AWO#o|v$!kWKOe?4yx zC2c=iteId9K7_JD9t}eQ<<3rPf^N*BzYoRJ1d(U);-kQe}qWQkdJzUOLdwEquRdwpUJY&eEKvX=&AX#ZFi7OirB-jxm z+==X6C^zAK4)m#J8Bix>Jez{9o0_a!K^Q1U0u<5VR0FS`enNwNQoP*+j=D}@VHL7v zep$FUZ4ezpv2S|TV-p&#INN<#eeGXn*r@FZNBR*Qv1r(wQrIlpo(aeffDAziD+Ns} zljRW^vRj$N=9B+e0C=7!-YJ-sCrp$EG3gKluerx*N(~7@JG|W08v?sbOmJ zcmMv*1MPiRVz7J?nPe4*Q&fTYLKc2#5w!p1j5rV-5Z(Ro_W71=Dr4dmyD8aMlX0sH zs(l5iGu9$aXe6VY@)Pkdv%t|zN3`@rXz30-NH7u7aUCKkG#@xggSncNSJ{!ZPI`-M zVF~t9Vg0}K0gQ2uNjui${KF4;1zxH65{Fb0L#mbNLlDW$!kKh$S#ke>;~ltmSv?y+f{>n^0QPF zf2t-y_sotwwE61=5|0DDi;Rm>skl{(4emTf^J%Wq&4SCtkL`n-si;#xGtGmfOs-X( zv6*|QNooFinYRh2J}(+oMvq=o-d*=~w&{@VFB?U?<*{J3rjIb`lW*c_n9}FIhPM_O z>!=?~j$HN5T2JVkg)hghSq}PscJ$Ri_Qa0`v|Pzj7}n|!UO$M#XZA`l3e4M;)d^#? z$o)ut*r_NBn{qdxj}A{Rm3iA&<-4X%CeXQz>RK>kRO?(5E5gmo>~8fBIrp}{7OLdp z5yeSuw}6OI@;0OJ>$iz)wmI+GTr$;QedPJ0k_~(Z-lVR68|%_5b@@Sja62xRRNL70 zIXUTqGYj|9sIL?eaHBi0f6tj_PdZXOgLBb##NYHnXk@jL=vr&2Ez63wQ zq`N?E-|jdv)`Hq~WLXjP`L)~a(*Zq5|Alc%gA;aGc{2IXo?RDhuK9wM+uc59qa1A< zzU__7xL2H6LM78hn8AcD(#OpVDZGfYF;>q7l`*gFR?dR-FP`I9A9fB52wnJdfV(0| zX}-;qoE#@76NFaOlrLeP%|=IecYLw@3kxMKBCD5fhd(r8x>^n^3zL-M*=S?@ch6|m zX$s`mM4c-lAM2Gey>c73G7I@78%O{2U(Q|{H`=hobkzr+*mihmj|ubHRc;Hpf!UgS?h@q{@OmaH4qG$e`y8lHeR zJ3ENIZ+guJqh=BpA4@k%&JpDX4YS%&5?5>mfc0S1u#xu!Tt0P(TO+mMPZxhrpK@5c zWKuuR(cwM09o1f$h>&Q&y2X`QGusy9lG>O$NsHmc^O0Fq@aBn)e}DOOvbeul6R{t&Gt`!GceY&%AtjArh&sM_{ zr|Y-<6rBNgP_^M^UDNPx%Iv*F*t=e*vX~OLZ8}N8jTv6tl>=@1()%g8jYVv$!|j%VenYeZ|9^VY^`YVG`p zwb9(B)mcmN4n9GnAIlNr5q)=k5G(z>4{2$|%%*f(-0#iB=JBt6Og_jj%_Cmra14i4 zWjmmbP(PNKIeivT)aDANR8`Kx(VHj=u-#O8A;Dsy)Vsd0*OWz@tq8eyf_q_C=XRrQ z@4zazD#4UT)X9Z+wy`N-Nz^pw2!Te&w8YK~jIYC*3-w`_zwF?SN(^Wd^)A?0b92#B zZfMkLT2A4AHAw2>J+AKESdu409QVx!GJvq+>1rftq6UXbpcFt@K?P$44TC&o|CFk@ z50E{@>IU?UG4BOdVIPh>ssn}eW*;q{k(1gb%Yx1k824`hDx7#SNuPRhvjCQi48}pY zPsdg$Wr&@G5u!qAHiYm1^>3%`R0bjkgPYcl8CD38eVEFkxC7C$7?)y7@xVxvy8heq zr3DQ7h|~uwRo<^LJh@7jlS1^Agrl)7xt{=*%BOZkgOcfImmkYjtq<~ImH3q$AWjzA z{^wN17{Wfh{V;uvUB|UfiLr!e+~vOFv(t_IW<#<9^pysOIMEW8m+uyRZI+20SWtt6 zOXlP=9>Mb~w85~xvhSZUhupz8XH$v?oXaPk-|N*DqO-T`4PqP{{Tjuw1{M!UPaGTL zrgm-j>YKyy77v7Gi?TJ?TF~(uR(ViAIv0!_b$5%>IW8c_br3>wz3RqsS;S%2l5|AxLLD z;+Nig&39^KSGkJ!Bfj5*mF6!~jkI5%vCE%bwF=?EBMP(HsRJJwtKKG6>MNnQ+4D|+ zzxJl?mX1l*pU1TgH?hc(T$!{gOiy>=03qw!$1*nNPVP%YocjivQOVL%;AG3f7_PMF zD=o%mjN6!}NU$d{(Mam0DW{HrgFpg^0;8>ath7mscA*hRNW^SUCqhir;Rh%Hfh_qB6$Mrqy2 z-Af9%pRF-5YH(P&bUZ`+j7ZrGoW+vujMJkp#1OQ5mK$r+>3y$9oumpnyCx8ha z`?iOue6cM{I>`*mj8Oc4>jKI6PJT)X^frlgRmN>!-=1xV5AWSpRuN`VE%f-l@)`shiesbo_d3A3{(11$XtqtLTZ-3uPbx9H3rn!*Iiz8aWMKMjaAtKf zjb&C8boco79foi}qMhXKP}t=}27x#IMI1K+z9%W>+hgVC!3R8nB^*GhmmK@tQLzq%%3CGDH59Feom;HU-qh zm&|Aigb}OiunCYewfNzhVA-^qEbBFs_ZQyn7LC?z+RJ=`(<+Mo#!?UJ^YRtXXyOco zlis|wa5(r2>`Dc*HBO_Iz(stEPMwOB_!Xt4uGXwi)N5d%Jw9LzU1V*fP{j0s;$vY= zQlp}dWvKW*w`ew+YN1Aq;2XeCn0LNoER;!|(Sx}Lj5?Ey$NV(HYydJl2II#mx}1(c zI0c{1pNe_+^f=M%s|zrbDchYBEP)g!18?YeC$xhg^WAaSKIqajW#d3AK87R1`)9nb z2l>2s1P`D@YGv{p&hJ%ZB1ZFXbdY!pF5Ac7#lOH6H}{=&YYMgMM36P>T87t{hNVTd zF6hHo;1!Q!w0rFgo)NqVRt`aHv6<@wW;o{(Rp^Gw0;l7KP4rCNdRjC>Q=y3@o zBmRI$?6!9SU!7+!Z}Nkd!tAzhWIg@o0^xy_9mACVi^aB$M$7aL6o!MYQl;-XC^Opo zr^<*c+fQxP7NiVpZ#oYmj~Oi#L6xx05~EU6(~Td+Df%?LG5$5@l7fJ7o1@PA8V7wVLsYZzwlEdHA3dpL( zWFqS#P!^S8K5%cU(s*T~G1!p83aj|!T3czLiq7sf^JZXy$(^dG`9R$I?WAkj%!+FLJvFeZBQ@p;uk#mG%m0(j zCTBt1abl7=8IJ?6a1?yV!X6D;bths&$;4#?4i1ziFAy;VhWi*Mp_MKhF%pk5c!fkT zV8l495j{w-#=;Iqea2vO-F}BZBr4JR4b|kpa2{=uS10n3S$lo|2+t&7NRf|v+?>Oj zm=)*NdDO$-v$(FnWcr;%1mzc5I&%fsLc+*pDJ3#u^%E0ugiWJp3B&v__j)V}U7C&$ z`F`Jn&9E8ywtDHX{=%})@v5%mZ-*Fq9CcWgJ$I$by{whooMIMaLSf(ughFCgD${20V4+k}TE z;m5iebq@TPn>m&Et+*r01wa=Ff#n!Q7LG>w7*k15fJDl(Uy$u8fE+?wj?xb!Q{_)h zXmvs(W^<(MN5ZhzIM{3i5yE4zp@m(3GK>(?W)^apsuyVRItef@V@gRv*7cjxB6$Pd z7;pTRf6`fH-(m)}`3v_tX$H)8CSR!cF?P~#E5Vxi-#OUi)JurP=xRa~J!wc%q)c^) z*$*Z2hkn!HIM3>e>*V>QgmJ}rVaJmU$x-%p*i?)8PJQe~-TK3DXD=?ip4m)8KQho6 zc_t7HNl1i;`%V>Tfx=CQ3V6h}0^;Y2D_l!*WK|lUp(0JipTlJDO%Dy6D3$8y$PEZO zPcc+mHcb-U%q1mwYu+r&#DFa!V75;x1Y!A&fe(;TlN>TbUyG#mEWk@LEgU!?;VEAaZ^bYGH>l$Vk-SM2 zc2U6~oPAEc6TFizn_-G9e6*})Rvm>>wj1@H9R1>$T_@A5Gs?#b5LPAHsJemiYqN2= z<5MPTk#`_E$g}gpFq!V2y7Y1iN`Z_cVGv{-+Z=Pctb{}EW~V-GTb_Oiwe{WO1)zjj zo%d2{Bq2uz3p&m;rlN+>#zf`)>SyF^)0K;q)s*^dxq#YC_2%p|5-274Chg>C$mXl( zhK^|;|B{H>jhC^Pko9Bu6id)$gvqQ|L;U>Y#BUmY+Hzx=f~ld&Z(maCrPibe?9FxC z0SP_Dp_U_s)!!X;xx85VOPU68zr?`JO>>2qw zaB6o`ViZa8;hW+^5kGis^X~G5m%e7T>78gk**HsEnfc@&2usJ5&`{P2QVfk0Pw=jf z6rYWD*>qeEK@B}BE{5#qC`uGdO5bJp^K*#C(4P>^p491Gjg1twN zEn!C^l;801OhBy?Zyb=S*#V3D0aQ8P~Y2y=yn+JMJ>LR$T7N{nJK zJIjmErGt8F?gPbpCJXTggNZImO0LwMWlz&e(&89cK!T$0ir+Khb4GAWWmqlH#V?;G zes@J_On@-V^|bP+5swI1-2qyKR%Vpr=t3)p=Buxl%}CELPN{>ww{_fQM0>yt^Y*G$ zFp0d{E2Ar59IU$J2vyz3%LpX#_sYTKEZSf}4CFL}Row&iwb1;HXQ?{yAXIkxBV(D% zYOTgq+8rk#FKl$!I8y`c5(S~MOq+ed{*;5~xB2&PX8LbB8ORZ*XuP(*Y2z|IsX#GA zNf(7U`t*1=z85;5KgC^6ns7k=esSF*OS^YBl9=Te#QZ#kM~+W;(b@QXG0)cp5pj@R z+Xm6+rOL&4ox6a5DUJ(G#js<&?IkaWeuASalFyyUDlxUGBp#f; zp|*+~bd1gW;ZV{svETJ^HzQ2$;?c2>BmAsGu7(_{uiG5au?WdBp&H_U;n)0dlgqAL zyZyS`4>cT4iaY}_^!pE1&N@T}x;ia^sX4Xe^I{|`-ff=cw)iE~8(JBeaak*J%R`yY=rF1Wb^TTAb_>yt%qlm3d zgvHBPWVEwTr41Gj%0WU3zYpa_ox3~9p)ray^L)fbK|a2lIiLTXy^dRU<+_fAiEr#W zOP&;y>FZT_l#K|?a(eIMr71pA9R*iibkDF_t0amQ$w#See9n&Tv9f8K0^#lly*&$3 zzfa>FDX+ElxPfeo7e_blUbhYMx&hhE$}sPSO&)IqN@U29$gwOLzWib;G@h@0tC&bc zI`*LwCF1cHFb8|b8l4521?rbF6huz9C?)W;J|=SG!YXO0I*P-d{RZ&IGOHbgk$%`Z zH$r8vFUl*)8I3}`{UJDoz8NKYBiS2$-K2~Jh{0&SI%_g_)Yd6(%*JgrHE0($Dm{lY zYx{b`6(+Hm{LNeIfR8tRhZDu!ictsTm}qYH2srLUyM*KoW|69S0uTGSmuJh%1kX>J z7TXp8;ggHrF&V4EsOpVY4~+r*K;nGfV}YKNuI3mkg0mD56#m@UuLk0$X}Sma6sfmS z7I*MV?rcYK$Jn)N_bPk8HJQ%d)GPhwlwXSqCBAh!F$7=WcdzVKZJ)X{FE`MCget_p z7>A%$7h(=OM4eT;zQ#G|qvgBKnHALIWd_P;!s;%=<{j;h!%P$DFf7K(^Z`AQ%lW;q z%kYZQ8?T2|Fprd;gXA2cK+o61vRa1yu|xE&(i%y)cY4Dor_ zAI3jzJC{bE&l76p2l;eZ*AG5MW}lqb#|2ZxOsvJ)M5^prl$+Rou+Y|l!>!VmOjaw3 zo*vO0yXJfZ#}4Qs`DLk`Bt&b_QXYMM0cOYS@Ccqg@LmhVo;=QV-|LaPO|K^j?kcCI zTzZ`JH#1-5!qj<9GCKr)E>LkV536W&+QHM9kJ8c40)ATxy)Q4OY#Eg{sQiW^-gCm- z*bE^~5ZyB#=jrIxL_AHYd{rwC}%+d zvRf8;mht&v{3S!ra%cKi2iPH_!XU=9*n(u+$veAQk5CWwZyf{38~MBE1G_!9m${!I z@EUyHgLXuH!BLko#7P91P@C&t*lauh3Hnw;A+u~}=o%9rFxm9K!KMbU51L`A@BVxo z2=xdT4SNX5mr%#Mf{4?l5&nQkMBD!aB6FTL@j9g$%^y-#5eU+)ko*#8Z+{e974WN4=MX>&;ADEKd>hS1YFOnrU*&;-DLnL2IkP)Xzw zE_+?%FT~q@!VF=&Ma-w|{1A?~3ur+5C*A!gBhJR|czt>7jY}oa83q3naKiwslch*J zdynp-o}C8}{y}2j=4aBzX^x_aTfErd3 zBz(HB9N%~F&$$j)sJjT==uQ&W78Xij-Y5J6M#K}01zCTJ^q*E?O=Du=mOE-vL%FNW-Kk)>Rlr2 zECJ9$e63B!;sNZVHz{S&%X;Os;y&lm&&WQ5Uu_=XiZWw%u`=RjQr5zFcmxgouXW6F z%;B+H(P~UC;iT&#Bpnfzk|N*%%uJ=w*0`%_KLFcoYG0T3avQI^*B&NZ@mcS^h}I`S1D&`MijJ=tMhgd- zUdG9#gSW*GO#u)IJk6s{wZChWmsJ2HbSlH=aauTJ2Endg4l}_pZSo2Ts6OJF;7Ux& z;arU;E4wrGt85N}GBoIH{6F~{eG=Y3@_e85F2FOUhH14N)8(CSe^t;H3{Bsn0tS^@ z8kA-Q*lfyaX%+=7Ia{PG=Nq6PDq8hq=M_Us?|IWD&8Alrr4Leo2Kf`5cMB=H z-Fci@0F=x;jgaeW;K{>O-G1qTVKMF zg{fthZtxpk4{|sK3Ik4p%mYsxy`;6mP$20BC6W@hRl1`yt5Gj+0MhcduehNYtV?zx zpAiIdg#vXb#?cmuqzz7TP?t;gpy=?bJHCWM5(@frLZ(OErb>^AYl)4PjT#)Mok5WN zkPD>7Vz+Z9ZX;or$!VF+B$5S`7VvLKLV~PP^4V}@W5-h+t?Ke)RtuKi%do-OY}WoX zP`@eIml%JDo@+z*HL1h#7I^r+TK;?+cX=yZ;NMuQ?d3Kzfp2ZHz>6zCnQD`$6q;`K z%B-9NX9QjCSrm(DNeUq6V&B{5n+{KlLiT8qmGi4ibzLI-hIN#F#aUGNrpfhA;2+&0 z!R|V17z2P#NMRNMYLbre4pTK(e_Yu=Gp+hC<(OSZ1E>pqDdI-KUELH*Ai}*$q!2f5M^)PR= znV%>*Vo)J%TfF{v7|!Lbh($$I`ljUZo3hfzk>93ZXLdNsQCHPg%Na{MYNCA0d$k{( z3bQ{v6I({_xqn-!Q;|q*V>DO9ux%1c$%9Guv?v?9_A;*0DSnv{sG(lLu!_=fwQD&p z9Huj)-Fohkk1Ty&&mZZ265M@~UrH(7&8={C;$%{>x2Xf+YU$=o;Pz&%YZfUPpwZ)+ z1k6Yrt%5x)9-dn4lkKWjyWR9Q=aCv$%Z||7Vs?3(u32;~5;G|n))v`-!y)L06Ki%M zM>(X$!__V>Pxe4Z&R-JJ&hwFsxjuH`QcqsHI#tEz_s-%BXDMw*MdkvRdW330lAIe^_)uM!6Kle7)>+V$K(^EYcI z#f6Y4v|$SVX$}xe&P$P%y^x6Bj<_#M=A>X8^(8I8xc0^77sM%~+Mq>9&tvCUt&I(A z%_6UVBEgBHnkJ;uIE@ljK-w!!Z+V@*cW`#Ksof=bwoF3gw?O5cHl#U$DT zek9Far@PkKxI^8zQ0|ydQxAIQ1hVed26ZkGCdHl?cDDE)oXSj^OKT$1+t>9`sC;DzPCKc= zQFJa+U*~*Sp|B~Ej_c**eU|po+I)9O+L`Bq0@GPO>*H@^>VqZego5wE9_SE zeOGTDP)eUZp{8#4IgeC&sDp{ltTBeSUm}7wOEZ7p-irMpn|@b<+HrI6O4-v%6q$NQ zONB%Ng2Q`=+42@2Mi;~GY|$$MilVCi|lIDV$AhyV-m$xz*dD?$ji1wUKr85j*F zr|5kVgkfgZWTxNl?pcNqm(+IaCtysLut6U(F8*~7O^H~{hAXt)p4IUBp^(s6k(H1n zv-AOw9q4+Nza~2*K0uS?>cu;Q3Q;KLR)w!~*M-`re^84}b}DYe{#9E^X)QLIH#4-B z`O49aPom^qxI+(>tLqX`&F|tr6Tbv&+D0so&j)BfDknLc`%VUoX5$Strr1rSI z;(k=))Um`%+jwr}z(elEHeR`5QkA~$Gq@~ENrtRV4&QV#-U!oZBu2|`;OIDXVq@OX zm)E8ALU3$5C}X-j;evl?x%fzO(y8(FfPW43$v%}PwsuBn`L}&-H;KiHXZebF!NMig ztqe;eH|m=5Fj#NgNtfp^kM4+|;1>R~)$eUzq2NNr1(|jScsP7kvz15l1iG+b+2dZc zLNs<)luoIja*-UhRZa&3?o-mNvF;)Nrh|yL+NL~uk%j7rk369Ku$Jaj%tNg`gf#4+M&@a;rTq+jYbLiy%1I{WgoTrs2E(J0<9W%~a z68lWh)9w*P(*-4yHu^e9JmQrc(Qtv89` zxR?*!jC1o}m#M4FMXoskYULY#jC<9Fl7A z@_gR8S6qPTwVx_yC^lW>sEmJVgePAX*>geww*Di3kxargcPt=i#NS(h)3{_Mq$h9qU z&1m$j2b>D#hT}n-uFn+zpy-8-c!$u|_|JUYQ_0<#@mJ+-Ql|YPyQxyl)FNq`9l7@H zx2VWn4;66CNpHGn4cgD0M)dLnm~izh+IUZ~aILqy1w5DtWQ2>NKVa0%+5&_?in?h% z&{P+Xv7cN&e@E9YbeMT69Eam*IFVdNTuO-@4_Foayo* z@Ya#CK-^ljz4ManuznohU)dgmV$`cD0jN{-QlCY6L)=r17qxvvw6O(H(R zbW3v7X!gR{l`8dSlLp((HAGvv>GyCUJ`dVZ(HE#>UYlwK; zc=!n9u-7PSt<{g8igY2#U(_YGj|{R=vYz=MwJ@sHA1sFttLU`I zSQs%9{hhM;F0^XEM6G0Fg}BiCQdOUvjQ}u5#rIKIK!1QstVeP(VM**8Bs4=1U7BU- z=cg$*bg=s!0t}O9`{d1g-B}fXw4qW12qN6PK9_KILe!32GjXqn-wswz8OuYzib9+i zReF9QZQA82u`_9@ZM_zKW3hhvBl%pqM&eNUpu%8z(#wm(osXdrh7@p^Y!9LKk$Hb5 z&DyDR4=XtBTXG{XPoAu^z84LNA)=(zF0#a%{zGm5nKWDcCu#OCCcA4o%WE8gJQ|ND z6rG$bdt@w6kSpy3t28xb9}zSYEpcLE@i&#FwF-bxS#hh)StX1eBP6wl@v5q*K(m3v zIyTjYB5tcG56K8-adK_5cZ%=f878`AN;8+$oT9a&{RN9ESV#ehd4AtG$QeY_mo{Zl3rwp*R!PhyW`&qUrSHjuzM46#mPLXnVRAQ4r9L6g7H;bay4}5|2uOT zB~sqKZgW=(FUXuV(mg0&%Ig<+k=DB>43UiuhQ=f>h?nT5w@c z#1~TOT(MK0=sRSoCM8{`su-wvGfTfY*YimHFZRwlDyn{e)B{K;EhR0WbVxVS-Q6Hv z1JYe8A>G{_lF}(kNe|tPba&s4Z=Ls?_nhDP{qDJI-F4TUzg!AyM)quG@9*<{o==1n zC@*U^Mru+KY^S}Ty;>-wLbhp}f%-)L(98hDLU7AcUnd^Cf(vZ%AT`lHe1FluBR;xu zgqsxoP2>I2WuurZ#jT}&;}En!=wv<SLJeY?=osgsbizYv1qgnvxzOG@R4O_;h=Qh z;4EEp(u6nf~__%BEvbJ^S+ec2g2Myaji8wt8C53eF@6amWct z>d6iP>MlUpT68Ex;*#yI8c#mj6=4;zzjGkSlCU`=Nl1^tEIThP6E)g&*~QRf604mDniZHQq3$XFtX?J@2G?VS+J522%RYtQNOrXss{VL*~Ejc4?dGVDbi7{ORgriw0QJn%t^!Y0Oye6M%kS9$uDRd
Ad5+7na_0NmqN(}vk0?bZH1eA$C#`tdZk31?k{|0b_f zf%Zj#Y^o~=v+dslwLc!5Tct7C)-k}D{YXy2f8rJ&t8|g`CtR!i3V5TL6E-~yY8i1h zk$%iMIAiW|Q}7b1s~T(Vn@kDSppun&AGTB=+tJZvLTevJv^J9P#s;^2WM}8Q-ge1I zV|<GsS=r?w?*Y)p)t935Lb%W~IvzJ*jB@|O;(f+Owr(PEUpMi%k8(W3*!~cugEcH&6{`%_I@2N3^3|(Yr z;QOtg5-m#jkCpJ(^!I#clZGe2pN4;9x=asD7r+v^R;qIqc-OZ}0mQ2dXgie69O9y_ z?YnJ%>(7>MEVAivl)qL)%A8Js<*FZXV*1Q{N{X!5=Y%6WJixz_8&?}JolklAzLCMt z0*A^W1=3}TVDq|2hZS7f1-@?jWY4=Gi&FwOF{e7qDEY~91C387zv4Nd6ONHGwgS*bg~|bhSKTnP5Y5`f{E}3&wO=A;pc&tdRYPrxA3w|4fZ{* z_FXM{n{e9op~Ln`0p-LOB>}cIr14td54H2YtZvbRY9-jN-sn|?Eo_w$?T7W(wqVAs zn@^4*$UkVMioZa!^={Zbsa9~iOsbzZNxwTxuw2?_E=sMe*fKIo9co1#t>Ece*U`&w zS+DSH)#PZU$&sC&8SWyW#!$K|o#U7=URs6HbRYS*!7dmKcGw=L1GrC8-ml!KXMM~q zx4AjB<$8~s7asr77k}`L_iLVJEHDvN zfcoSqv+onMbr$&CcURiUtJtO>&=xKa&wDp$036}r2PX1!`h3F{hQg;Mn0!h}QQqd@ z%L_B9-<;`(*{cT$PG^ZoqCnT|6E0@KOa$k9`6d*m74YKI3irhr{G|xULO_>nj5i3B zo<=>CDn3qgG*8^ZBfm6)b`(vWc@8+b|M;LJmIG&%EX$qE!>}$$vs6bQjv_f)zXx>R zfBf1E507SXhP0&Y+Sc*o{`RXU&GZAer`s}=tI*X#sWXd0gS_XnuTR(geeIcRmS;ezA||FuZ!aesS`#>{BN0rI(mK10NNJPI7N0{S1ow??1OQYj=F}H*0@rQ?_R(IZZ8fa+={A@h% zrci)g6EMb4vbs&*Ka>M*kzZBuReWT-JJ6cMphVn&AMxu=6wqL2cmS<{Z!Ecn(|RGY z+$}9bX7W{t094%zDAUa-1CtES9>2dq{s!S*>gof>oB6klTPW-UDYoNEI;lP1du~^9!nGNPk(Bynvxa^doMQaaAl?|R|jk^Q3# z^>?>gq*|3=96tXf@>2eWc3CY=q7Q9DrwjzyA8R?Z@-rAVYlXlnE0DJ=$;q~ z1uT7l>KBlrxNQi2Z8emOy1vn&MwR}E{YnUaTPg8r-k}UGVNq$RX(|H_pnHm%fz;U&T2+;vNv2Uz*q=_)}<`m zZVvJhMYdwS)uRO}QM}LCF6ePbu>qVKy-#7qwpTVh1SBS$H3nfKa3;D2M&eJ}8t$f7 z=$kDs$FhI=$%@jg(?aE?pf zFrVT8=j9sUmDyRd17Q4*p=RX z*k&GWb6zZ_S)vq?Qrk>2>zsDl&$YJa8#qvBb!$U+6)r8#4sq#1XSVpxIo=IEa9!ug zN4Q9Ed1mVr1ZVI;+PeG4V(A!Xf7prpFO;_AZz!$UD<>riQ=Hz?sF9AxCT738uogL0 z|4kQmC~Yx37Vfi83fPbDe++mr^S)^5|bc5~aT^`ANyL= zF*5I{)XU6beXq+crfSG#^($`$=Vq=Mu5)2v*B&rltkZ8-V;2;BtBec2BD!nlP<3w( zWD2s8B53p1EhBpsG*l~MpVmW>^IvJM1R+EjdQ3_GQgi(VEnN|2ZX;-HK;X_cJ07@g zkbXT&M7LTABqz0rAA!0h#tE0A5$52H%FwUhc(moTprWg$_DSFo=>vN(QDSR69n(qn zSU3X0ngSr_G_|U0ym|*jP(DC&BRU7GZx9kvu&Y>57Q#O#hr^nH{lG^G1CcckKIgNR zliyfVKDo6j1wRAH3pAWZdBEn91sAtD<^i2UA5ow$W@`rQNqdt=#IU01O7+E5Pb`Pj zM3KY+_zhmQcECclHlQJ^xx4W=Cu4H}YYlkL)p`FQtM*HJy%FkMAf<89etaUz-*@8E z61%xtr{LimwxME-GI6UlDm`VEsgrWGcc+tt2xsu50(^~4m2U*?9}5zRkpq)Q5sYWr=H;R_wHJvSh$LoJXUx%? zEZ3eGK6w!l6uX(8bBKhP{%I}qr}OFr*uV83xq6HI5qwq3Jx<-;Ejo@w#(i4AUl&(Y z#?urdpa@I#!vhQDr`p*KKmwUu>SH_E(ELOh6hE9}(2^-gq8Ebtc|%f4aq?}*-vQMt&gog=LbI+*A2N8Ks)D%P znRnKX`^EdcajL>)sQeiEs0Ks(d1F6y*N2F0l+t7TJ}A_?%50KiXQ;Lfv#mMUIB8%*A8(YR5&4=rzRnmr zrpqu{4V`zOPpUJ|bdyl7(UOGNLOG!Xh=paF=465VoEkAnz1p?leGPrt@JlM)=?ojk z6^GMs>eQ5yBaeJb;HKV7jlS8Qfp5@+1h}|nX?}6%RT06EVu09uc|@Klbapf0rU2Q9 z1S?S4Ll?yJDWIm5qKWojlX5?(93|^Brkbu=*&% z^OX;0qFqNa9!^1`zd+R!E>1 z87e#(vlo@PBvwPe64i}@O^=)ghL}ic2VymU1&Y$eFKD0kH5Z@mIIv2Hef013EX-k- zX&7*EoTZ>0?G};7n|btM8sSorM)T9U*yu+~`%zx$!6~fuz8W6p!zZ4q$_d!>KI