From a6085d0a0a1d0278b52a196f55ca56058a3d64d7 Mon Sep 17 00:00:00 2001 From: BigFranklin1 Date: Sat, 17 Sep 2022 17:23:17 -0400 Subject: [PATCH 1/2] finish naive scan --- src/main.cpp | 6 ++-- stream_compaction/common.cu | 12 ++++++++ stream_compaction/common.h | 2 ++ stream_compaction/cpu.cu | 39 +++++++++++++++++++++++-- stream_compaction/naive.cu | 57 +++++++++++++++++++++++++++++++++++-- 5 files changed, 107 insertions(+), 9 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 896ac2b..4716eb3 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 << 16; // 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]; @@ -51,7 +51,7 @@ int main(int argc, char* argv[]) { printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(SIZE, b, c); /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan @@ -64,7 +64,7 @@ int main(int argc, char* argv[]) { printDesc("naive scan, non-power-of-two"); StreamCompaction::Naive::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..ba9a3a8 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -24,6 +24,11 @@ namespace StreamCompaction { */ __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] ? 1 : 0; } /** @@ -33,6 +38,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]) { + odata[indices[index]] = idata[index]; + } } } diff --git a/stream_compaction/common.h b/stream_compaction/common.h index d2c1fed..b76a9c9 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 BLOCK_SIZE 256 + /** * Check for CUDA errors; print and exit if there was a problem. */ diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..818d30b 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -20,9 +20,19 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + odata[0] = idata[0]; + for (size_t i = 1; i < n; i++) { + odata[i] = odata[i - 1] + idata[i]; + } timer().endCpuTimer(); } - + void scanNoTimer(int n, int* odata, const int* idata) { + // TODO + odata[0] = idata[0]; + for (size_t i = 1; i < n; i++) { + odata[i] = odata[i - 1] + idata[i]; + } + } /** * CPU stream compaction without using the scan function. * @@ -31,8 +41,16 @@ namespace StreamCompaction { int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int j = 0; + for (size_t i = 0; i < n; i++) { + if (idata[i] != 0) { + odata[j] = idata[i]; + j++; + } + } + timer().endCpuTimer(); - return -1; + return j; } /** @@ -41,10 +59,25 @@ namespace StreamCompaction { * @returns the number of elements remaining after compaction. */ int compactWithScan(int n, int *odata, const int *idata) { + int* boolFlag = new int[n]; // temporary array + int* scanRes = new int[n]; timer().startCpuTimer(); // TODO + for (size_t i = 0; i < n; i++) + { + boolFlag[i] = idata[i] == 0 ? 0 : 1; + } + scanNoTimer(n, scanRes, boolFlag); // odata: scan result + + + for (size_t i = 0; i < n; i++) { + if (boolFlag[i] == 0) continue; + odata[scanRes[i] - 1] = idata[i]; + } + timer().endCpuTimer(); - return -1; + return scanRes[n - 1]; + } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..09a4235 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -12,14 +12,65 @@ namespace StreamCompaction { return timer; } // TODO: __global__ + __global__ void kernScan(const int* idata, int* odata, int n, int d) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + + int pow2dminus1 = std::pow(2, d - 1); + if (index >= pow2dminus1) { + odata[index] = idata[index] + idata[index - pow2dminus1]; + } + else { + odata[index] = idata[index]; + } + } + //__global__ void scan(float* g_odata, float* g_idata, int n) { + // extern __shared__ float temp[]; // allocated on invocation + // int thid = threadIdx.x; + // int pout = 0, pin = 1; + // temp[pout * n + thid] = (thid > 0) ? g_idata[thid - 1] : 0; + // __syncthreads(); + // for (int offset = 1; offset < n; offset *= 2) { + // pout = 1 - pout; + // // swap double buffer indices + // pin = 1 - pout; + // if (thid >= offset) + // temp[pout * n + thid] += temp[pin * n + thid - offset]; + // else + // temp[pout * n + thid] = temp[pin * n + thid]; + // __syncthreads(); + // } + // g_odata[thid] = temp[pout * n + thid]; // write output + //} /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ - void scan(int n, int *odata, const int *idata) { + void scan(int n, int* odata, const int* idata) { + dim3 blocksPerGrid((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + + int* dev_in; + int* dev_out; + cudaMalloc((void**)&dev_in, n * sizeof(int)); + cudaMalloc((void**)&dev_out, n * sizeof(int)); + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + timer().startGpuTimer(); - // TODO + + for (int d = 1; d <= ilog2ceil(n); d++) { + kernScan <<< blocksPerGrid, BLOCK_SIZE >>> (dev_in, dev_out, n, d); + //std::swap(dev_in, dev_out); + cudaMemcpy(dev_in, dev_out, sizeof(int) * n, cudaMemcpyDeviceToDevice); + } + // works fine without doing the exclusive shift here as mentioned in GOU Gem Book timer().endGpuTimer(); + + // send the data to host + cudaMemcpy(odata, dev_in, sizeof(int) * (n), cudaMemcpyDeviceToHost); + cudaFree(dev_in); + cudaFree(dev_out); } } -} +} \ No newline at end of file From c1a3b528a4131039bf96b5912ac12c4bdaeba958 Mon Sep 17 00:00:00 2001 From: BigFranklin1 Date: Sun, 18 Sep 2022 22:08:25 -0400 Subject: [PATCH 2/2] finished work --- README.md | 108 ++++++++++++++++++++-- img/table1.png | Bin 0 -> 33163 bytes img/table2.png | Bin 0 -> 30396 bytes src/main.cpp | 57 +++++++----- src/testing_helpers.hpp | 9 +- stream_compaction/common.h | 2 +- stream_compaction/cpu.cu | 12 +-- stream_compaction/efficient.cu | 163 ++++++++++++++++++++++++++++++++- stream_compaction/naive.cu | 8 +- stream_compaction/thrust.cu | 12 ++- 10 files changed, 325 insertions(+), 46 deletions(-) create mode 100644 img/table1.png create mode 100644 img/table2.png diff --git a/README.md b/README.md index 0e38ddb..b058605 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,108 @@ 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) +* Yilin Liu + * [LinkedIn](https://www.linkedin.com/in/yilin-liu-9538ba1a5/) + * [Personal website](https://www.yilin.games) +* Tested on personal laptop: + - Windows 10, Intel(R) Core(TM), i7-10750H CPU @ 2.60GHz 2.59 GHz, RTX 2070 Max-Q 8GB -### (TODO: Your README) +## Features +This project uses CUDA to implement and improve a number of parallelized scan and stream compaction algorithms. Following features have been implemented: -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) +- Scan: calculate the prefix sum (using arbitrary operator) of an array + - CPU scan with/without simulating parallelized scan + + - GPU naive scan + + - GPU work-efficient scan + + - GPU scan using `thrust::exclusive_scan` + +- Stream compaction: remove elements that unmeet specific conditions from an array, and keep the rest compact in memory + + - CPU stream compaction with/without CPU scan + + - GPU stream compaction using the work-efficient scan + + + + +## Reflection + +* Based on array size of 33554432 (2^25), I optimized each method by selecting the block size which could provide best performance. + + | Methods | Optimized Block Size | + | ------------- | ------------- | + | Naive Scan | 128 | + | Work Efficient | 512 | + | Thrust | 512 | + + +* Based on the optimized block size, I compare all of these GPU Scan implementations (Naive, Work-Efficient, and Thrust) to the serial CPU version of scan. + + |![image](./img/table1.png)| + |:--:| + | *Comparision of Scan Methods vs Time* | +* From the figure above, we can see that there is no significant differences between CPU, Naive and Work Efficient method while the array size is less than 15M, although Work Efficient method is slightly faster than the other two. However, as the array size increases, work efficient method is much faster while other two continue to grow linearly. On the other side, the Thrust method has always been the fastest method and take almost constant time as array increases. +* The naive and work efficient method are slow since GPU has to read from global memory, which is very costly. The reason why Thrust library is super fast could probably be that it saves copy/paste cost among kernels. + + |![image](./img/table2.png)| + |:--:| + | *Comparision of Compaction Methods vs Time* | +* For the stream compaction algorithm, the GPU method surpassed all CPU mthods. The stream compaction method using scan is the slowest one since it has to do more operations from global memory. + + + + +## Example Output for Array Size of 2^25 +``` + +**************** +** SCAN TESTS ** +**************** + [ 24 27 7 15 47 21 25 25 4 30 41 18 28 ... 11 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 49.3806ms (std::chrono Measured) +==== cpu scan, non-power-of-two ==== + elapsed time: 48.9979ms (std::chrono Measured) + passed +==== naive scan, power-of-two ==== + elapsed time: 52.6029ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== + elapsed time: 50.8384ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 21.571ms (CUDA Measured) + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 22.2183ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 1.02707ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 1.06106ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 2 2 3 3 2 2 0 1 3 0 2 2 0 ... 2 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 74.2421ms (std::chrono Measured) + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 74.3851ms (std::chrono Measured) + passed +==== cpu compact with scan ==== + elapsed time: 166.162ms (std::chrono Measured) + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 26.1186ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 26.5175ms (CUDA Measured) +``` diff --git a/img/table1.png b/img/table1.png new file mode 100644 index 0000000000000000000000000000000000000000..a7f93183cf65d9d81c03837f447761f83324050a GIT binary patch literal 33163 zcmc$`WmuJ6*Dg#6BBdf-3d&L%q+2PGjzyOsUD6$jfOJWxGz%6bwP=JzN_Pt^y1QXt z=zaU}eBZl|{cmr72nR8*F~=NpjB}jhyatNBuT9@XH;XyPxh-KYsfF&GgWODmATAA`Z5u6I+8xMD~|^bD}_Uox>!o91sH-%n%ItZ0WjcqR=Ut`8F~Ttb+ZERm4XXvg%6*<~_FUK+J%l?5+e9?#|& zw1svgF0ymfsg2s#oc$V<6AC#?G44rv!4;5RTbp!tu!?WzPTU0?pad8U-ZLWcKNavg z8aK8r_wMT6H0w-zyXNIPUO?dof$2Y)@ErwZF*3i%p zv46VqGyWyJ>A+Kk<5n_jc?E@t+86Emu16a$>L*b#aPLNK9-lALDiF8*kf8Dql?bDJ zRx~Qb6dQy~$xqDrP}lP`t5&`B-PD)!qfPoy^Q#Rb=j5(#qtfiB6M_mbZD6b8ubAA? zk%X5b9ev=mYs0F#1kD0DvdNM8ckpSY?-OxS3ZlV$Y?|laF%XP-+nq;7K7R_<&o?XKA zZ~jwVk$0m<*>ogy6yFgu&^LQuToxZ))Y*oOw^4bKb8U63IjN9^1bYelLxBs$1teb-vvb#5u?ZNqPKwV8?~-_PAkCN0)+4=tH- zPui=8Nqc&TFFMD}nImueH>1xtxqhhYpVW>TeIr`it0!K_HDc~M{Z6^w7)jVF95|=e zPt30`@~7{)&z&NouOvDq*r?vkV{kbwTjCSubpU8$BEG(OcTa#)VwHITf6 zcTIB6>v2}5MS(i|86A(3e4p1{B)^~lyE4DSNku7yI78sLWp{S2*BN=S9QEwOw*NPU)$_w6w>N z3*f@1C&@fgN*w*-e?jsd;c8HqJxVdWO@b6omN~7rTjW)*_kkgVV5C5ureZ@~cHZ>rVPP!aafIJ4sr{aB=1AH5KDkLb^Q+#6~C;9TQWirg38(R&5WF(U!t54>Ho7q6yK-VNsB$ z5WiNJ+=dM^ZtO1ein5vuqqt{L-mlAe>t7eJ9GOH~*HkrplKKE!zAc-%?a@_xKkFl- z1?9o*0$NhmrQU9FC~zs~I1q&gv5KX$XHM*dFRiEx(gPrJv2)t>?5D z(=(Ht`zF7CE-#R>(3Blh&WVR_YaPmpZE>NN-)-I^a)r%iTUVkpxE>mP6C0O6Y>+gt zou-@KSx{wNR%C15mDcKM7DjKq6T+;I?*59iv)e{_$rjb1uB`0$E)5SjO9XCC?Cgw*zm-ki8YL%=WH1GnQX z0@CjWYZ}PkWo5?$9qj5$~)c zOsa44_E&J`JJmNm;BQ&~7->L=z|u&!HPBuxKQVg9j4L(+idylAU&mHM#r?YY?(P27 zNw!hX1KOLtynh(o^=BJ&AHp7P2N|4U2Ome#oz&9K#!IK~ImbqyRK+TgD4u=Oq?ru@a$>Wi)^ zXU5PBRu;v#JNF-X(S~NQV9+NrSM06FNh-JPl1^4y=j18n1opOD{>*+o(<7VwzMz&c zqgt0%K`sz^RvI?LUkk6J*HY2ewlpJ=;;sKKy0{xYNh%;@0XbpzQsI%G4&H?+yDt5^ z4e~Hs&8r^7C21xv;Q#Cpw+|-?jRIc^p5Q&`H1f+^HE(}ib0DzbbRHV|>GBO4~JE>jY&t4I>()X_Q7sl@6 zHgZ@+CdIasf7gp6o1vWuC=9~eSx=NN0Lf(QD4QF$}C3JC4CSqldjYEnh6q%aK*vw1YDFa{f>&`*GW?^0X871EKujRe)u6($N-JT~nY+0Fj zwkYDfQ@UEGS@`$G|1XoVb)`=}wG%xGHJxu%rX-x#P2CF__qI+;!c4KICfDn7`NAC= zM*f~&x2hn9Z2f}${;~JY2igY0+nnWe9x(=IAWfR z{}|Umr< zt8lzEy^83FQaHTcrMiBgbG0vY9i>rWAqhb6AV`=*`ku7+>?DAA7WEg83nrKd+CnKZ za4C5EtUdS1ssa4#dxh&SVimJL^`RL@Sxt>tCza2su+Mp>)@6>Z@OnS>np{U!Y_W!o zYGDbkcXe@U4<}HSPZjh?zfzrcvV`sT@GYhX3y(;RFj1rm5X3`p+8OkUg_E1bSI^dq zh`88t1|=mz$-kX%H>A=6=VE10|7NcUj)_kzpRWR)A4!x{PQl!9GoL6k9g>-TUL*zJ zb@5d9y#N+491b7)ZY;OfpW#2nu*-=LIPf=E$Jv)nLs*&sKdQ3 zX8+Methlbl^x3PU$~gzv%wq)@xa{Q};0I6Ypz7bnr^Xg-eSlU?K8TQUWi7vmkk%e} zH0!2S%_e6*7JO&6;bM__{u%Y*{gFdg-(9?%wJn+--_c5~_oo+ZCMYqabFjb5E zj&5V{;%LKF`3g^77cs-QaF&KL%M97^bD)tHz2J!L+Ud)M^_+N}?k&DEXyh6@Tpta( zwys{y{;MxHbo14`wL}4K%ALQ-nPQ+V8)GW}j$B_1m97&>& zV@w1nemnK`{!({dS($R2%vBd~yRHt4=A9A#^w?^;BIY9!TO_QLq9NbKOE~vci)7V1B*x zXmfI*0`>8Oo18%5$H&M>$A3I!$I7Zko@hU+?c_VEq}|B<>k{A=6bv!&C?65E*vp@N zkdg|zw`V-BRsbYk@#spc=>&8fQqdqHX2NFPz3wJ&w_Q-E#1b?3a{YC(Tsp z`o&eCkOgV%tMi6y&A1mN!^ecy0u(0so(850bDNd3gyONO5cd5(5j+K8!F2tcF$R+A z5<&hU%g0_i1|CyYHV4NI*8-3p;Tdt-i29{#d=J$n&W=#5Q-PvFFP8VaTbpi~Ef3`|q)wc2;GudoE_F`t+<75js-~ghqhOJW4 zQLBBsskW!iKuk<5d*Y{v=k;#qL(?{5Ipe+ikJsN9E>pRGH|`}|8oKMuWgWXZ;eCdI zjx~fO6~plDe5)p`uzHP-b&B0_V=UM8^fhr9z@~k%gkhyxGr52N`F`Gmwtg(sBu zD9VlEgf?YmsZpe$fMRw7$N9~A04U|YoF0T#) z*(RXRHbpz5UnqCzmA>6MMLTX2vEgC&L2&tP*L%bU=3t@ge zLf7uqbADtd&3-u9jKAsKp$>W<4vR6CwejyDo@ChsPRj%qb@ZEB1~6BtbaZsN0CO=^ zsLeMRbA!f^))3K6)j-xT9n6-oE%9~(>Pry(`UbrKMMK5x_nK+qIwS?Dmm9`ImDajv zt8!j%xn`VcC;)Cy%V`%*6g9g|bnV7ZS-tZEtJ&fAsw(aDK8MXzSK&G?iv*CuLPC+x z-HBXyqn9fZLietz?xKlm-u#?*LbxjWdYjNy$(IuSKlQunU_h#4c|YJ9h7sge_E14y z{;?lviC)vEJ6MDau9^e_c9LCZ5sN8|POgx+KM5Q%V;rc{5{o(4MKRGT2 zz&Z!s}drM&wolwiB1~9OcZX2xCh)AgxeBT_oRNPU^ zsVap`OYZI3ec&reL|8+(|Ko3_Sv6sr;QY5a08?M_5xevzXk!{c#DHgSd8=kdmbqGVzF< z+o|#b=W8gM9I5#7PK7~ryU~fA6dylKS1gb^kS?ZYZ}Y7TcUaLC1`cgOGCp_6D8cbAu3F~9;86KDm|r(J z{Yc)h9*BZFBdUum?W3K*K@)#iE*6^8l;f8u!p(!A>f>eSmE`JPh^%jparp3;=3n5=;pN}T4beCPPEmxS`n zB3`v%9XY-2lV9gb6Dj{J4WZHIE181?k`I>x9u$aOd@#32Yr(NJ9&p}a(%i`3F>ggu z2O8Ab#A>3tuP=36&!$&7>X@QzGfo`5g;Xan9W|t9gH1q|DIrlM>V~>?TFs@_Jf*6>72`=H=&KWs}8gMEM^X zaH{>hu=IQ&GVB=_<;Sq2*cMSD^4W+x#_Z%+$JJb;ojaNjmAc2wye|{xLuvKnEma-= z3*|NeHSJ!UXKdnlF>+WqRD0%ZH}SZ3j`oxDR;1_Dm#+yl(@2NZO6%a0{S-7UXC|+; zT5<-Q48sm3T>bHt)g1RJB|1U};wVBRpkGUc8)G~cH_i3850+w|G7 z9xFLsxn51`g!{^)r4BvaS1Z@aAogk@lU!V)we;AvK{1^*9*!X?2Z(ge5*4=?87Fvf z%b;_MT917D(eCM-dQp}ZfzdTa)qGld z@c%^w81eoelLAP!eBkg6GuZMbjYQ;w(zI`z!(^7lN3N`+!$dSne6uFZ8D#?V zR#+k;Rq7eZ@*vlf%2Nl-07TCvLzPC}W{UfWC%%7N5C%~^aB6boBE@R3V$E!ik~a+HP@JuyQ#$NK@yIGo!q z_At+ijS1iHV$LlXoK5#b>Pm4-()ymQkOp}>^Yj_uQuEwxe3r@EE_Lni+y?b_W-0G3 zNhuGl#I)c4g&LA~HuwY~9v9c_TvBYc7ikTK!u|| zgl=|Ac4oLMpsf9AZ#gpHX&eU^lqHl~o;IjdaW%XHGxC%)Or{E;L~+kXg=Ndp)}>3& zWR3uxRYB7hKFjDW+bA^41|0SZHN~Q@-9rMmYXsQP#gh&jH#rL=gs3N^FG0@q$qyTmY4@0}xLEPk7|%8$H{eCwC;peo1` z9I)k{w7uojFQh08n9|zfk=X^Bla$`?swvNsm_f5JIweXlT@z{Ck8;nMe~-38Y3Xj?!J(k*(MzJ~!O;gg0|A!kU7L1RAmXAyw969he~fr;@G#boctq_=!4nHg;j44z5H`~wqb zNq|5;BZ83b;GR}k$0Sb5J{B%J{UUy8?3TNF{y=Fsx4~y7VJUmj?GhnXRCLcEsHyt; z>O2p?*-WAH0M#{A?O+ZgM%)jO6~0e(>mjV>SwsYpcJpd<9tD=l*5K11xhS_LhI>^> z*^KLcBEGyH)$LMvo`_&llnu@qneSZZd7T8|T#|eh#0);b1V04;p$PqO&@X+%+lXkq z){4R<#$AFyASEXqoy%VaEO~B(Am;m2rUU|FDJPdsSZVYucW`B;3Idcf6QHmpzo?v5 zjxQP3I|LK};=#nPe0_cAcO}Ag|9ff0h1tqPi2p&(_oXL|w3SMWHP+!y3mWO(OxF7J zvew@9(_1;QAlT}RC*YMM5s=3LG}bW{GI?y!s^W9#Lv`_Z4#k#AtDxXNnFj2BP^xI_ z!4c8Pd&y{aw|E2A*1_S%dk2yam@0?ysur?tDX)! z3AY&Mzt)PzY)Iv&k5zpE^4273#=l`clI3 z>7hN!XrdQG2C1TLU;&;^v8d6L++4Qr($OTEIE2Qcn}8JKO-G|kah#@gUG&k$I88Yk z?2UG^H|#H{=}?8DfTIx!r_Z7tT-aUj?yUN^b9g*lB`^f;kQrR=qqR zSTfiK2nKE2j%U;vdO zc!7FL#UAbCb(cmmLAl*{jTNjpoSeuc=!#$Xr@vOb_f~t7)Rd0OR!1^Ib7_ z;yJ+M<$?26Y?Z=&GS|_568&4dWDCs*KJa!3Y?IM#kGuc^^`OzRJBuzHh<&z(K8e2`xK-$H2nihk1qnWrr*eZXH+{7?m`=M1oh zR0-rTiJ46CT$9CJe8!2zoYCx7*}{t^emqO5eMD!x9~7Nl=XL)}EkSb?8I;??dK z1g8C)57<=b8^ZXzh!o)K%!&#bnm@7wl^D6j@n|+%6kM<-f32DjS{x4!ralTVHN#c{ z*2KH)k#dt=)+u(C32+cDM-%zn_}GA1h2-Afd%zC8z&KVo+`-Guv}^`|f(MYdgmh`P ztHIRjCfbeoAi!;P3o?b|3=7rLaaufsZ)%#@@QX085P0#9C7t$v4}jJrOZ97%d8;y= z+tOU_lg84XurIcCrLSjNI2}_5Rx_;Y(w#b4z^#bR)StX>WUW2%kh4*6?~aMnY890J zZ8VguHFnE5)>RZ1`2<>8I_I7=R^W)?#a#D!0{eUZ9Ix+$K=6TtVQBsu`^2+(DEC%i z9v8IqE)dPMZavnw`_7?q6>`4CTbLo}OZi@b8iRh(8j9zSu+)PE(ACeI+<^e0o?^a= zuO6Wz18()Z(&Y>ETYa&twh3JBba&A6Vkp(+Zg}cFd)eg2DGv*nBnJ*LrLs{80B|A(bQlQnE~=z8SfE|4F{>x^6Jkfuvf!Tf6dN!AeboGt3Cw}cNv|a-2ux`X|DuhsNq$+B zV?L3fOwKr9dhn3>(o)U-p?nHa_%a4VQ4^wbb%c)Sf3QfFw~Q?t`ob$q>Y$7Rukw){ zTInhT>n#M8-7+~Wg#@3XVCG#_R4e20K*zvX6&-!1NWNCv?;CyFn(}_SxiXH+*$R&d zQq;<@l^P1qHNkNioHH?LaDJ04&}fSA$z@tUY@9HOW=3Q+X;j;rN{1F|*EsfbjB3%O zIM)6{Eh!u}^4=yC3e28mL*sd$bd+((xu>-Vc)8DI z!!$3IH!CPr4evaI?vZQasC+l-HUan}ge^=If^OW`-5ba1b2HfDg1)ReUMsBG3S?EA zasz9TyYe56l^ED>mXC@LqpshkMI^JMultZ8SI+W7WgJP5zDD_s43}{Lqpkxu9hJb9 zsl67sErKbbY(m1AeuD#S6lr?r^_W@l)k&MsJtadww<71lA7`K3ieeZw2@Q(59ngd} z+nxU~W*KWA9d6%`c8xm_x(<$gDLS{K{@f?RDXp!ChmTmHA=An*RCeQ_)GNU!yqGJf zB1s3kubDT63>r#;{0j)az*T(84Z8@}#!6@aC4n-cL8IQCXBrq*lvz%4!c3F(4s=*c zo6h>fJRe;N-M`z1=XD~Y#F&eD`(Quc`1VijWi7FFzFRDV3Zdkbmu|G+5)3hc3kn@3 zY!Q|y`{m6O=X-wVFPkeCcb#@nf!}EZn*q?(LPwGD@MFIlynTCaK#G~brJ8uF>1_;l zIe;n2wbY6M7VjBX%;s!Qj>vYsn6b2ow=a%W`hs?9nt*{s|I*>Qeud=m-tciv-tvjk z(8VRBIIp_Wn!|xa{Y}Oad1m=~Ws8N+vpTHV*+H>O?qt{3oGwOD4gm4j4+u(Ht^qzU zZT9-=c#c>8<%!;%wST%bI#%S^sHuHp1((Vkpih&cYDi5|5(L#{zL|NtNzOd7`rS~+ z)x_Dp`0RTs%w)oqDthEGrir%q0v(5oVe!v)fVrq~%o6lC3vZrF)2VX}_$DLIck9GT zPAaUkA;{bEHiHeHWxzmZyV`E<1)60#7f~Bym%&Q_aVi)o!D-EZDi*P5$K!_P4>jr~ z8!WVj{3djP0tZTYj;K{jbG#LCuUEv3J9u_q(s^V$++b{}-pKQFmq$BQmcyM3G-sYU zsjkuivM3X_;Zf+_`+Y|XoKN#hAMllT66}rWA)o(c8$qc+dtQf7-=hE@IHv@XPLX4( zStFzIDz|t2= z*?p;32cZ9U1pyQ#)XcW`G@&!qi~0sRb~qv~_OJ%J-Y&{8L`mmcwsGv1Q3)Ly$I*F` z0(B3OgL?X+$(=f0v}&Mf{+j%FcIp=3)>@EC~lcUhdJFHgD3LS^1~(dT9_xqm>Om z2H~WndPoYdf3n~yT*KGIj){1topQ`PL+<_9=|8bPSU&%*9)dTE-|Pthlx!#YVYj;W zO2dpPeoGVH{xcEM!;mPngUwY+4VYZl#M3l~SJ_mm@klov+26X!ra+JT#>s<20WCKX zb%e)zTEGeYP7*Fq=1Pvsm=vE+F+W3IRK}w6vU?7S^d7kr@2q7=`X6pR>Mi~8k5lC{ zHBCCJBUL)DQ}uy+vSP#g85f4yl=yYD0ezuV&z4(Cl#eztsQV_;0P;s6yHV`|<0#N| zw7`_rA*tvEI9WS$0MfH_-tpa)1YCdk*1iXVu}#y~{_VAO4Yv|y8YaER!5B#KVc+kP z1B(32Mv?(0M1tNP9(0uLQSk~tvn2iUEyue{5{!B-FV@@OzkW*FYuE2GT5$g`;Q9Ii zm#rtOxHRs@bka$OF(-OfSjCfjpFgEi7)XRK>_;f(-|tx403yrb*9jt?4l((tuion`=oVb#6a*@g4;7DT3pFzZH$GW`&a8$Ze&6k(LzFcA>`^z+ zHX6S7Xd4Mf=9L&9#8VI4eBqWYqUZzWsG@>p%LEqSZOCl7%*JVmzELqI>L@D>66MM7 z4hAw23sW>rH6gMup#%_0?(2+k{{E=!z3CioIeB<-NJ?|S4eJd zc+jg9??2-edW~Hc-C+!mmTOlr9)cSJ?YveN&bak(sl)3-YMB^&2X>5(`m1M_Wa8a+ zs{IRh=r&J~2)HhTzOe3D&Za9H6RDIz7;WrS%Bv+11$*1m&$#5UHexF%2Ujk6D9x^3 zJuQUHx#6zy49f&L=7Er*QqWRhWJC&JkL{T{G1}uFKVIK^8dpb-VER*S6@~-5l{#h< zB6Y~$+W%a4A9V8Ak0z$Hm!%|P=f^|e9tbFs18lMzQ0w{7oJihY)Ew$7$ymA(ffxN1Y&3CATVV}{vwuzRyy%1?y_+kM6g6Y-~ zIh=(>04?~D9ng)Rn46mfepNSaF&_OL@`sRF$?NZLTB8tj9H?5<;RTbD0l15=tP`9D zHv47&?`#PatwN5fd^997Z~FSBt5`gXGV&};hoD*SS$rE)MALj`+#<68nRc$}59c_{ z#N77|oRM$o=}(R~v)7(n!i#kQB(h4Ouh#cgE0G{61&MJCAZ%!>S9PND%=W)U37(Be z)yQTuzY$E2nnAq!nopxP^ zmKn=!h+!ir|3m@+ES;ty&uR1+X3FWs1n@9aYR~uyxEi85YaH$^>)ulaD@sm^v(~Ba zV#io7t3<4QbZdNn;g1)Xv&sqJ9Z^H(bf%gl=oQJ|1H2{O+8Vh z9#xGZPz;HA$kyg^+6~V&x#?-^3T*?ce-qbOd;$0isB{34F$i#eKFbiw(3<(QV>0Vn z1v5Tofk{>}w|HoOvuLlzSQ2NoFWro*I@*fewq7-lLVD0i9&U4n&u}W`=EGeT zIQS@m?xjjurxOqb0Q8U5vc-f=TkSTyD|Wt|4`Sv7n)!Qw>q%bOe9nvMfAE>>#oJWp z@FZQ?$8}|*09XB$7v)SBZtvH7l2&G)kf_yP8b0${@-hat}@K=mW!J;5cm}`>m#{`w3LFUJ+@1K-Ts= zT^s^5M0AWY7cIE$6o1hSK#;~h^0oC67~zCeLN2J1J1#Hz0ggGF(L5|0A1aioFsn{! zw&&TR%(1ndR+$YmJ=GC99XP(P2W_&snd^o?(XkL*e}TD$PD3a^J)zQXa-lhq26=rm zF#z^pF`#&`X=acY=Xyr!7%w0+-*p#@>xO)zbMiO_!!g=`$@TBEoxND%8Y(weO#oy? zJmk1AQUv$EYv;|O89~>peG{d_UlOvU4~D)m z3EJ^0@DvK){ge1eLfdYH5ZcajzHxLQxyIi^cv(6qF2Ndg>_EyOFRT{%f#~JLP#=DG zTcCp(K>?r>Ryb~|9V(G%^O+QQ$Zo-J`GnchpW9TSuVq)zIepHlKC}UE-_h}bQ{>{v zvHW4*v1wZYiS53|9Kesw0Q^{kJy4yBUB9QPZwgI{B!-z(c&hM8I%xfgR^DU|^O3%q z?RqId4TX1kkRy~;vsE1gyj~`fuUZt7_=)@a8 z$IAT+F7L9j4W7}e<0e3zqZl-YJAYoiDIa>*S5yLoI3;_$GLKO}2o7fh1omZOigfanj)+1Im7nDCPQ7)&0 zVE2hMC;z~PH_u^yu}4XYdhviikAe%X>YZA#M1U+-RXK3(F@be_z5djCGYlwI$mV|kKnNajtOe}#OC zM-g86Mv@wOQpYsCNhGIW%OC#^a+246Y~*6n{7wejcL_hB45j2Bm{osf<~rv~WCDnl z5w>tVq&qOIKkiX9h1W?;Wwb=y>Du7`7JwG#kLA#Rey~c-e8_+N`3cSO$d0=U{gGH= znUWzd$93b2Be=NX#F_jR=;dyO6k3}sz(GGOQC2`+FLIy(Sot~c!vwvNz|^gUnObB! zerhq_cM#__(h#l5k@Ou)+8n?zWN|_ zgOT~4#o?;aKPaew&Z(9>%OzwsXYmAQ?@ujMR(IfVfs}=ryv?Dn@nN*ZqgN6Vw$_WnKVER97bv0Mfxh zm+n<$58pQb8JyYkl^p*rTJ+bi?LZleGzI3|08!sz1ih*KbApsd-${$D?$D^!U51xF z2X3uplFg;Px+lQwKS>PBucnPEyQgdoy59535ngnqh9(>c$VF6?yHx3Dp3i1dxp0{@ zdXljV2T5y+XHdBY9W-yZ!&!^J4X9i*@sqtyznnBH)+jT1f->o1Yb-PPxrCJ8x%dkl z1tv$E>^+Wjpa_~QyT=Xb;5c6&!(Q*zz*3XDlOyx%0MoXZ7cT(#ulK>-hd>?&dqhfV z-06&o^4wc!cRlS4M~>_pnGT=Q+uRa8TkS>66Cc#|MtWX|z_^XB0TkkO8te6KZEY== z?f2*=qf^)_Qu2>tD32ZrsDRbX0pMvB&$T7J!|DM#B5#$o_Yjqj5GX>|MQkUw?FEdk z6{!O)u>PCyc&VNUXcgoa;RR>*0cjpYHJ6cS$K4J#5_l~Ju8zS=^jSY2)Zn0)2r2smh^FD zfvY?XqrV$YQn?CDYl7+U%f#IsM-qFmqx5UzYcBP;-klz4YRmn)Y{LHbbN?I3iVaB5 z4+!bJ0F?mBEJ6q0K}oEuy+h`6o>jM0ngNmQK%AIas!qe`J%!CryTI*A0B)D~M8$~c z%#WQb=nqrs`k7z&QSym01Uc%z?6Hy@UqvUFL>B>)JvLzc&TQgN8tG{Pv@yOG-f_Z) z2s;z=>!1Ur3=xsfIYL;Z1-fE556sA{aCPX7K93p#e%CK{rb|-BYWHAUL;HyWj88cF z>H0H<5HId$KbNl%jia_4OrFEyAS)1kwY4Nsjv6X}(7n2TZL~gRFb}*aj*yA`G5)Jz zaxMpeEqSjRzU0k+&Q2n>-&c@%dawMgOj)NYASH~?1gE|C{E0`kyO85-rOY>u6jAi` zf}RO0<@j5Y-N!YGOv_q@i@6#!V9jUOaD*MKKtA3f2 zS>Y6ceo)n_6bYs&Y{yeZ6c4f7ko7GWU%HM?I`U(%1)l&TgIOSMZ!$dKkFBOppAfH- z0a0Z-r{zRD{IUN(0Lroqpv;qyU8SPhtM~Pun|>RWf-`Q1W~!Pv;KAxIsE3|+WK3M- zt4;RnJY|$_0mOkgDcj#vd}x~1j?sxyYCg=U2!Ubdec)(sVD?~rBeVQXAm6<4ALk$b zy@FYz?9Ef+MQ%Spe>^Y+vlHVHQa=G+WA6a)5y~GjtGxd$r`ptciwMg6u7#I+AGyGD zpo>?aOsHNu1+@%*L=FShN~TMc^Mv8z$?SZksAYc1AKlXq*J1Sy1xTcLRwhA+|4FiOU? zS5{F$ptNLwWd^v&;HG=&sgx2|1lO;wz^W1TMf9fTMKpDpapdQYZP})8c5^v}x-+-E zBFZ+8qq6a3v=%c!5@HznPQ+W8y)iDUY9+3~ED)eN0U)FZA}ZJS0qaoas8;)p85JY7 z!}0rMfru>c?_^N}bX&Re!qMX!W%16L)w?n0++G$if53sn ztgGpAN6V{HM|m>n9*#+tER)$CpF5>~TKveHH%W~QitPj5UE8OAVnVFTx_T!yXpD~Y zO*%9?LWNQQ*$!A0jxSEnL?rhlV(IZ^V}7Qk<&%d#d-fc_v?rY_U+@oXyRr6D6Nw^eBJU&pE(vhTqQ^>m4zo&S$U8vt}@jUIt>rR^hh-3o(yr0s;o#1ki(d zYFl@~$hQ+51Wjw9C8?rxm%DbFd2CMksT%J0@2M;!QWHbpdq=(X4pe!J_S>JgExhTu z7kDA=HGBjd)z6ajzui!*fSv)M`Cu-JC-2C09FzK1q5xSS+lUxcCXc}RGT%Xvu3{Cl z(LP+7{yxs{6+7vC$}Ps(A__l_SJ}{9&J#7Ok)o35Yiq5`GV95 zr4}Z0_olhpD+So2L%>G_E9@2@eI{^1_%GYAdXd}hP*dojIcao3?HV6jes4OzfKAQo zYCqK!m{lR{c~{eem>!yvaZ3;;GK8Q1>kJY4>nj0ej$$qXt~wfvg13}ruX;Kb-Z0Zg zSG?_>d##-*=X|oTq6NmoaT>^dro-#Z$p>gX$W?FTzd*bLj{sWPkIwgu=YIIwU!5-d zgJgZ1I`M3G(aP{wve1pyX3sn1>DKqu04+(AjB@|}HbT2D=}o24CvKwQW%}A706c4f zH~n~5Gwv`=xVOKsjs*Xzx)Bpc1!)*YH|f)G=(pT8175AZJS_4eFjUHx{w8#}DyJwF z9`t7q^gJ3(FKmM^0C^kxmjjh@%jca{J@-IT4hh8r63LlpX~%e&i;xLVEt*isJe4nN zfcSUp+xEA9&;!)(dAYf7zG}N|mX~>6y511#e|A9N?=1@uD|z#WL}JX*?+Roc-NUD5 zdo#Ry>8&He0r=@r(kUSq7Rm{>6tw@X!?I;Jfn=&V5Eopa{JRH~QyGs`Z{#+efv_VQ z%gbxUGhpC%5rq%th zgANsL*HVFg2vwzh{AmQERs=3%6_%;yT!sUyPaq!RWy3UQgUVva5cFQRMLqAzHfv0h zS8-_2pP;+#qy$z2KVs^OCOpsk-@T_mO0vQ(wzI;+OKDA~GViBhsTW>Cdz~SU4)}vNr+5CnnZ8M?r!s&RyR-xI zQSUR$h0SdAVTA61*BAR$anim|ertFUlm&IY=&?-wmc+U*O!+3FJXaWDv`ghK^?X3B z>%o1cbnGQ^2-Xw>XRm-&Z>DjquH4aOmX9zstOvcxY5XG{X|KYe$mVdDT;~(!h&&Cq zVx#O%Y~Zva2sG{#aUoxQj5Wj5QFTVK&B`*#zUv~deGo;nf`I%w+<9Yd-c`RWIeYm& zgvheZM?d|rj4AQzt;P;UWy$-Hrs^AoVKpxs8p>XfG)%3779`KDouvHE-l<*h)Ig45 zrl(40dAi^%xw6`TD&huxEa3GwX01wKwqL!JU@lr;NM=>_ z@#gPD054vAmWCC##gPLls}G5JANze41HQ7(q{|nmj1TCMVFbiMx^XflS;*VUuf+x` zW|>D@>MMv-YUu_Zrv(JjVf0e*IWj1w7@#G|vVvt+@$tkzjogr0=XGK?O3@_+K!zc< zbO6fwdJozIS01)Mk-J08fM4p8r$!$iW1a5C;~at0*_2Y!`c$Xr^HBbM*2hHfsr&m@ z>PP#;+Vf27*W~%;>Cc=uBO78fffFEUWOX=rZ6(aU7j6g@NHGmXdMHrhSK%bWnOyvC)b0OSDO>%FD)_6 zUQaM>ls_FoSYVOn&{MC6dQj_W`}y?Y<2>8rbhq`JrMzEm`VbC0tBu;>}xJ%pS2uHg+qi{&* z+-l`{D>K>=YMZQ_?iRzOat$L{e&@un8tjtp zB9`y8`)50Xz7Qbz)7)-)!@o`MnAf{ZN8)+}E)&J@Kw;2~EffG%%3smykP7a$`Asz> z@27JxGzjSMS2U zk*Iqbw&17Xo>amZ6&ZWsNBsTLiepgvE525A`v&Y9u_{(jbi4V!56(Y&4hZzt^^O_a@cRh`g=%=4(e65JnCB?LMwXV?&G>20r|>#mEs0Ra zYv*w04)nTYg6uvsCuFM0EE7>5B&-k>9x>^k@Q9#cUl?^6UQnm-0hhW!e}40RHDt$U z)Ah@VLg?x5fakl?McIhKt}IvGdeD!rbbJFky+3Q=W-l>kST@oc1*zV1cbAX5EC-Ok zs={Qu|CCjgQ_;5uBAoYeIL1^89~l3lwE7#<^0W76rp5fC)yEU>V^25~kJUfbdF2HS z9o?I7{?8bI=0r||-KPJP=V=zP-^1p`H9!nwYGzOZs9mTGbO|lVTy_*Yt~}ti=a!35 z-UA<;xhzwHJd0`w>L5V)Zb)ES3_v9$wNBQz0HnZx#i3(8?fj=Oy^$vGcomzs^zYIx8V#jqw*j2*|OgbZ(w< zH~?yt>!1*7SpEv2N(wTzoGL2gp_yE}W$0PBb*vK6%kw;I zxxc4juyydtYlWM#8GVf-x}|E&E4*W(1RAJ0DU>T{_!Oj;n%lY~w0N91q zBIoW0)6s(pKP%TVG2z+xL?k}tD+JXD* z2^xyy&bUN86)V8YB~0PQN<{RVy}iHK#E|rUtE5M{1Y41(I2=lA=zJZO05tMRyYseh zP+d$jlV0|(UKCH-6Zh5LJl@v3?PYJkr<``HSr9S8zZ#je+S4HVM97ym<-`W zqAs7Zr^3>9W^L3l;kF5cjM-JtkUro7cRH@JQy|F}3QdG$q^>Bd>j)?*S!?`LQy4-; zcb|If!^c<1=1pm@JIveJY=1_dczsJ7z&6WNfj<3**k}4Sk^@uTx#Cj8SnTB>4vbG0@@sT{8EU=9U@b6rtC#flp&%vkfNaFE z$Gf{o#ij|~gml(+sW#zyR0?5nQbzgeyy%gy0Tj*MW2%PM6wD!=^)t+pnX*7P<>a`2 zi$bD@aS_8f$oA57ViFCjM9s@fk+SG-fin+iBORS+w4JzA70F^VGyMLkG(4v+zl3Gd z1=g?Yxdtm^vNgX_J#q~HfQm`_toRt)HW+ z_UlDiIYgg2IRPh>2=o_E+lM-&gzF+``f8C_$bl=H|Tts__1DQVQVVkULq90C`Bkk3ZwH5=WeW7wG2>mH1WI$2 zR>D2fE&gyIL$CSZqJWjO(2QeTqDut^%rvvcaN*gr_!VqXU&H#sIlMPMSkJ#Z&C?Wa zT{+R1*(kktngD1ECex%eWv3$moeg%6e-@5g=;|X?_rfZ%c0tT9rh!C`gAj0P z2$Ysp`1Ksb2I3Z_LH{r>mbs660QdI3MeX!SLIGOK(+QQeM7}^tVoS_EorzP_!nfmb zL3M|D01;34ou|(uA>oJ@sRh(0Q^kIhpUoYb+pi?>wcRwMzTk`_!V@NE@bity+*EB% zMpRLwV!Dx_Y?tJ;h&zGW^vi1C5SG)en%A*m2Jl9}%)OYkR`WDC6Vxzi)TAWYHn0bV zykFvh&_Wm`^@1Z^< zaaN5dwR-;t?0{T*Q!pAsq|_Bif2NexrQ|ZRv)G+T%{+b~HhKXft2{V46-IO@P^taH zzx~ry0^k|V$C5B*qF%XCE?S7pcNx@SogQ9BrK#>=`9YZT;tC8?x7lkSdudk-}(t zSYq);dmLo&wQ@ZJBn%q-8B1Ys9dA#TMBo7)ws;p}46r$$rPkg3k}y`ZqM^NEBD3!K z3Ge~~KX7?s_prd2vhbQ1)R(C(?kCwS3STBIitM_grmG{CUu|iX z#kW%|81>WmzM2pd6q;bq$iNuuDbHBhHW(HEa|`+}viDouVgj+*8MZIFr6^r>#dlfp z`1A}=3|pbK)I48A+TZ?pXo2|ntF$QVgxsgG>g&Q`23v;&AelhVtJXBF5%0ro;CxBP zZt2?OJNuawBv{R?mZLMi;f-beGKoy)?JpBK)J#WfLTn(c?d9KPDU{}uGF>gQ7L#1$ zPl=J3{msrM+T{bOjYlo=jUIuH!1@1?%KfJ+8$u!gQGYGu>wq@gf=#7`&Bh{&uD=JI z1?#7ax}r2UU2-w-=$+h6sS2O`hwmj!~@iJRLh3oy$F=~fHG+Yy65@@fLH)^nY> z_8$pBlhO{vsB75Fi0`|h!(Gs@2K$~rF4D6$P; zl_~Sx(Q?CVngG_{7aRP+y`#4d2Nho%LmUEMolr{BCDO?5X*7O6_2CtQ+hssQg3)A=i=%psjL9<$=P?KI{*Au`!ox_-O`y;Zq}{L zlokULD2%4!>|V#D4Eq&+j)wFDDQNSX{xy?kjWaK_&4eaFzg+TyIET=>RgUtoiPoa{ zoWt@XF+Z=jF0sbD%7T;RfSqSl@(qG+92PcrQVd~yh_ScGrdn?USKrAuC{z!>Q8Yy$ z{b!aksdJ$Y_jUIXbiH2Q9u^*sU>2q#3Jn)(CZe7I}0^4w^3A3gkYb zG1bx%t|9r`tBogQ*EisqA~|NfY0WT|@dk2}({T+~oB%gqjhUiKN~G@RR1Ym&QPyKks;2%M<6l)mL*eM5UJ z9)L~&04R)EEx&q^qg(G=?Z6?QZXlE2uJ)EqEDjwqENU%KfhdJrprVF>lsB+N=^soY z5HW0K&PsX_XhVp`8+1H~;`31U!P22r%;M}rhHhl2)3;=a1ZkbO#Zfuj?pN;UGA@u; z`=*m1i$d8Cx5)I6WgX4?d1yT36=r*f2l*nx;Nidu2~tB`q+!t|9Q=?*tjyTZ(n{&S zMI9dTj$Th4@##=M9&()Fp1D}Fy5IHUEB|7OY}#46vo!7LiRx7;We;RVn-9`$qnGjh z8p@)Q{Myk~YgEod5^m}fWULqt$LZgi3V#rL#>V6hhYKl>;wIb=wk0ouw*(qC$5zvT{%ZH#yf!l&7(kjAtPI9=Mzekf6~!uO#P#IYug@#2 z2kvlYP|97pWs6fjw7CP_p3lct>O#h8p4jZFCR!gdcN99rd9?IrwCF&m18tiir!-(A zkts6~G|=^|pL96BOnrfkGJn{RcJ@EZ)h}_BZ9`to$(5756cw25b->tQ1H-ym$xWRLE>g(O{ePPcv7Gm2c$;s zr^9cPT-vR$r^@%YnC;6qyub7 ze+U<;UCK46wi39}_>&00Cda3XSDnXk@3j!EMY(glF`z6xUZ_I*$d+SSentl%Oh01A z6o~PI9UE`H8*~wtLF!H?|&8yw?X&3G{$wX{=SN_FV;;;TaNr2-%fHzJ5-?3pR^YQ{#tKbPw z-6A6j8jX|$0TptUbB|C0L$PlTEt$yx9$hrXT{mw9+bQtw=zCfZPU}#eCz{IXF8Rju zaa!8^)8*LgIP3tF*tXbET9J)AdAf-hMaBs^VIr2h-#39#EgR_nBbNj=V1pET&h2`J zBP;K3b)i#ofc{be0HS_G0(vO<`x`m%$bK0&jSfykvsVnKwxTHr%f|fe2vVLpbw_I9 zo|frS2`WG|6J!D2+E6;cq$g!BJguk|%);8C#(*icu>e4pKoWt?HCyEwm>~0MTt0)1 zReo{+9z^^*^DEVKGge^9i?W9K0daZ;k+iw(=SU(l9Q<30WapvIiSn0%@;1>A-l2D4 zRQXF(VuVv148xR}Lm_+|qQ*++Dg@tU1XPVpI(ulodH1tF_wi!#67UjuM?%+6bpGL! z1QQ*}v&f-I3GSAQi|ZyC6+j4%7HbPt=J)}!^Gl_!5yQH$=-wU3gHaNF6_%6Q>8`5n zQyZ!kFZ;{n;Nu_=n|?%0>^=-dAziIouw#DZob~ttjV!1R#soWs_McGH z0q9pQjY{D&&=ww=E7X?rcPG(EtI7k8wZPlPYNI!YLMA%nCt{21d@Yr016DgsRoNXJ z{K(;24Y){Eh*3dnELMX+*W;3rWXF%Hg+*o{w%Zf!IkiQ^ z#ENh_P&(uGxrZ5YyBFUmOmTJD?|od_+DpPrS;0ts&)XwI4?RI3Z>i)p(H=j{i~Q9| zuj2M%bmJ67;V6PbUV>hIv1_{Xjg*;vrQJ<5IPakcd*)!XzGF+{beuUB>}<2*yhu&%B$rdoBt@(X{yDA!zb-2HHyOx z4&*T^n>kEpSE4sk`|t$@@5AKt&!#hPI{rYCwxoEceQ(HmVb{SD(tHTivft(j7$#hL<`iD$#mD-Nr~9%Yo{~ zB|_AH;SY6hZOE_TpmR)Ck#=no5B~D)IQEX`kJ8ppUD0uxD0oV*+gbO0Ltg@Kq=LU zo9?PUC|A-ZT^l->D}1j*3QAcfuIBWz;M~ZLUDOh$l`}P-8xG=WutUaK6yb!lgi>@} zwyU|6CM8N&$QvX|`4;B8hk>`(7;IfUyY31YU^hpH83eiBK8uovhlY~U8JP{NT~vFS z@Gf_8)OE9CMpw%SU<;Ra&*un-YcqEy;4~tE)II0JhugD<)}{X?!8EEU5~Q(CG`?^Y zjZ4=(e=`?e*fVDUEB~Y_zlRS53?f7aJ*z}(y!vleWm0g?mp1qi7)7I7O99ph@? zfW*TFCvJr%Cf}}ck4DeoJcy*O{NnF@i$n)x{Xc_EqoWQ#38;CAKXTH3{**rKldV!i z7U+fP=Z$g9h={ou=dTFLHXN#n?qJdAHjH0@e)uyg=!WOze^IooENiAMk1-oQS>t0( z2645%Zm405@_+f?xkTsCN_*N?sLh}UE!KU*?gp_t$L6U?TP7~7?+QUu;Ee z8)>qyClyT&Sva*uu!$o)xAU+4_NmG%3?~`{%4Hs80FbgUONNc^}_1V%gFfa3c69d zqma@bA0tid;k-FH?uff1kJgLnTI-BtPc|~8;JN~v9W`OuMO?+7jkYay#as!AP?0OW zSENq5r7I~57Y8$<4K4hnN>!~l^i#kUc?Om%;kYZ*Gk=8#7s7VgX{gqV(c=Bna7M9h zi$B8Q6AIW$wMLN~BzD01@Luk~kyOrTVIl}V_U|J2HyvV7CGwuqCh`Xdc_e1(5(rrH zLYsz;(}=0_43S$Eh9MSIsB&Q)NW#wjH`PcQvGf?mr6(pW8M}nw++F)Q(x{kNq^1Jk8E=^uC1yE_4ScO$O;LzbG}~(5~IRMe_H-gk*LwtxkqIg zRLEhw>}><;2U4fTi-!Rh!IrCUzW*aQ20A!odgX`={bSC`?otItRDC1fGecvG*3RIb zHt~A8IyEuaj03sP9@0f@^!~g^gL_l&NdqtUH~rjhLz6Dz6W-nOdo}Wmpa2Klx*1Vd zr_`)=<3*+AH4ZnzagIwTc_#~q*PH}zuc7AQ9eK5n*82{)fq%-xp9MMgy{{J*-f_T1 zUNVDa@96@5^jHZYR~Vsc87Z!Y9ASBJ(rWV|?F;qNQ?`Dqt5Ogn>`i3>NLs;-o z3p1i&*sjEI^5Ch_UW{8%zanuz-u)uU*i0HL$6NVbS9g+IpF1Y0ut{e9rajX4bvY^g zg&mlHSDYv*CP=E9eOi<>= z?#(YKZO?euQ_ooZlDl~Jlc` z<_-im)}GOtq-K7kN#^qVndSctVscDf0sXm8p|m>HYuSGcbyR?P_dErNRrwM_gqFQg zc$+$6H*h1OvelBBlpHOVTQS@2!7tDEzY+FxlhpPZ-_?d@6&rxw%;$%Myy`&W`mp{r za*!Ju{h5+MsbpR2aDO62ux=D7LP3u96Bp&ht z>9Jce6B7h* zBdrU?d5+xBpR+vqQsNPW2A8)XfS zPI?xpF(#qT0sNc0ql@TXf zYQnWr$Bm-VXTn7e;IUbM-Gs<%q0)=}Z(2pVIwQIas}5S-=TV!2Xg@`(==xT*(onf&Z)q^(|oPlANe8RBhPlwrYg;(>t{ z2C9Cu4s_hS_NI0a$xyeJ)Y#|phMuz_Xr>e8bxD?safp^c4mR7f;$H7Y7KX?j$;00R zJ%90u;X)T3M2vpvy4&>8rsS8tRMm}BJ+{EYZr8Xs` zg6Mr~Eg!Q^cvJ4|I#)?9F(`+5w}Ja`$H^vf>x7W4SPGj_LCiwR7Jgpg$fZRh^9B+(e-)92Cv^6qv^BIY!hMj{ z-gdA@Sme3{Aw&OExX9^Wi8UDj`dvo`ZTi`3EcVNBdo)K84Lk~`n(2^k$2g&o831uPbIoaBHwO)ubOGJ^pMOHJPmoe{%Wc0rV zsepgYWif2@C1Wt^uUo-2`I@y;ZhJ5zyeUDkJ`&d%oa&)KA6!regA;hOb(vYIkFa_` z(V%Ock4Q4>;ha&#JlmJ?>TI5U=%(sHaK6@!!QX*NF&EgQX#CmXO9pr$urx(=lJk4e zq41~r3#lY2Y?}7vE7(K2PrmwibT5dUpIJxILMk}?p4B!-ZzBEL3RIPF2P4;pj`qLb z@lu%q-m_5`VV+*_&O6y7P%%8NQtHj*;NSCneUPSFYSHSd-_xh69zPH58{weP0<)Q>~KSlS|YPyB-i)_3s{b7;%l0ZqJs zQ!iJtXf+*`XTX`9#y|&V!pq2bG;mV$@kgQWrM>N`cQ7=#N@3v8J!5BeIe^-1%cej1 zR$3q*+&TIFX5&1w%0)xD6HT9!kO%v>AW?P96x6szv+Mh7j;iCrqM#>0We_Q*eMVRR zIgg>ibn-J#O_EiHQ!8D%hmezk{X?5k;w9U(TE1rCidlrAmKZn*V?2qunAH6VGBR@C z@-OHezcP)YyPpm=0lYRFhOog1KUZrV+5)lKP*&@5M#M%$pQ(81Y%w z9*JC(!~Oi)j8XuoYlmm|vi;MSo5Z<2ceUvM<(psn*<7i}hRIT$raU?XuBtawm==m+|%cyA5}(TH0u|TCU!T5{GkR65FUHv)ODm z0=Vbt#)TCVv^f~oy86lbg@p5yo>m6yB8YjT?sTX>kX-#mrLW6raDZ(ERt4}$AXaB9 zL*<5$O+OLfoAf|NHofBaj`F|PU;oX(a;VV)5qrfj&sH<}{?a;^y<{(Y(@HS|3_dN=m(SYTOYgKj?C!U``p5bbAtW zsbK*kfJ0((02(pD&w!$z)4cN%t_UosgQteR3w(+kpg!g;{Cq@;e1>0}gl&K2ADPIO z*M7*4IV`XW$Bdwcp2AqahQdpHwNX~f!fchHnTkO^ccd{^l{Ncpl`;l%(U~VlCd>vF3%+pj&8lD>^CF6$8CQzCavl1#3TjLcB6`g zX9I$H``2l<5gv)lN<(vNu5g!m_5&qrP9Bmt31^#O^rpOTlMFBTd+UMXzY>|lExa8g z@)D?rinb@D^9-poy{u%_1^Exh=#}(^ z2OP!TH$R`xM8*ot*M&sawGAt+dZ{c?DOyYuU%0Mp?or=pz`m?`YD|hY*8$a;^oc}m z?ja`HkM+&kyJP~|{;N%KPA}L^y=UCC?y=n!czI=vd8Ow6j7SiM;?y&y&E^#4*Bv%5 zbJFN8C6-v5TkLfIk&7@DlU}8^>pr8u;fx5JvL%|j7rKB_FM_NZH%PW-DM_#HSm(&J zv{?uv7g=A`0b70^3Evg?oay{ye=eH}25Qk?2;Qdtn22%lkO+1rLYcm*iQnw`c?2Jx zGj!DdI!6s#;(r+_1cQ^*GAHt;FB_Hii7Zq>lkQCM3~X*u(kObF&AM_5BOGYo>91u~ z@(Tq^JHyucf+)aRceQcJPO{aa5<>}NB0YBW1m2v#kLG!R0+9KU57(MS#beP{V2B>)6$XEf{RBH!tY#2y#YJ}d4=_}B7OoJ8QH z7|y$rz0Ha`(&{>t?M{DBqFGQm3zFMd4p)S=9+LNT7XmGjO|mtri1n zYqsU!VrYsVPwG_o0nNr0D!)T8x|~~Jge{WOfzX+zGSWz5R|mH^rvQxQU~PfKgNcwe zC3YY-r&AHwX&KxW-_N$biF;SG1+=Cs@m~PjU!ZaTZFz>jwSz?8PeM@cus8*&(VNx|${Uy)m`sYwqpnL3S+vGk~-OZRX=Q&B`y|NFmvAPT>+(-LP zyG?eitt-9u3zdQcKX0r+*M98dk082r6}9K;F^?NM`3WrGFiX`R@@-pj8S7cG&DLB4 zt4tzW;PJ_c>E+S#qa6MjO=6G*5qHWmwtImADwpXF;-cMyi|lWqBr&I)xiVt%-S(Ze zx$Kw9tqX#sLe^0?bHc~*WsE;X3OfSG%E)&#fwKvAOCz{G1rBMbQkoJx^zGgkM6L~z zZW=@FJDK8!5khaXeU1}1C`!^jDd};>;$;>G$~t%lB~I^~-C3J{8Bja3Uvf*ZZq=5Sh`Nt(NAjc7V~ z6XEz$5vkd|12;+vp3p-oe;6t~B(ky{QP<9PWVt}c1<8$znFa4EQYkjjnoiL(qcI|? z3$Jlt9%~!fzW$xQ&LdD;)?<)x9`MFd6){aepxsqySs#3gdJxY9lgHbMs$CdAiF?4L z{6rQJfsqs=X6fO;B3^aCAJ=r*_{Ln7J?6U{}1#uggmTUdJXCFRxCstdvw% zUYdty&7%*o7xQy6A9}K~sJwvh&rTPK=a#d` z>*CV#v@0ob@pK$r=dbri0sOyc%bGs|<_EC@<$4ICMT~n|_s$Q?+>a&_@H;PWRBd8h z%f6O>3k=KvFz^L99V&f&eadxKI)gsvHGKBXnWIItdnWhy0HrRNo0@?kHxQHZ(b$(E z?o;c_Y&E^#b+@*!Q}{ih&aaQFO~-_LBlQ8 z_4Akd*nluKBNqq1J37wbu8$LiDp8vidfQ4qVqEbS!)|{ zQJ|28_-ILMzub;+MiVm=L#x)=s+o~dj7C_tbqYL!hRrQo5Fkyck=o;tBc}s|R1FiD z({qC%0A9D+y8iC8^*^$(e>ed5Q0!ejJ;7-6sl2ZFfP==4idVIekOncguEv)Mtn;cR z3h8)(M-8rjA`(DAKv1swEW?a@F;io{zW%$;8yRE<=lgvY(7Kr@D#C#`4e+l7**CL{ z6w;FKnUIGO^UL}mq8w}a-_dZETr0Li*AD@-T5bT5i+E}e6f61u@xfpmKLAG-xfWb; z9=4^7)tTxY$rm>XXosxxO-?wJ+~j*|Gu{(%80Yk!b4bC5&C$Rsm3;tSDaz#mftBhe z{1~A6_-G9B=K%o6&r+G9$6;Ung!GofijLc&=>H!Zv2#F-Hgeu7FF&6O;6_Gh|MNOl z;Ku-!8#zD^7YnI;Zch7*Mqnlb4vlasS4Hjj@GbBC42?bDxI6}Ei*o=H#sS(C z2w3-OT#u4O1`X2|*q_QpL*I^PiXhWaQGH)5%0}g1biW|vJ^yXEvMDMiMmtH@w(EYm z9PkDopLXMF(Q$tU!J>$+%nl6Bmx!WJZgFuio2jN=-`Kb=FUAXW$*>y5uSYa6Fff+9xl|2am#ow*VFMhb3ss6@)%gJ9nn%YPXI{W` z>Mh%s@76vq2`HSk7~-UEe*%VYiVRBrm@kl;m<_wXO{IBUQxo!gXy0l>i@EH#G1>l;+a4a2BwwyP2DU95PvyZYBNZENMSY<{ZZS0+z8}dCSU!BkPcP0p8TGM4v7sVz+m3us(iPR@SZN$d2XjLR|;*O$8oz*DD_a z>HxU_S~aZtWUY)^i|Lgqvu>cntku5&_+EpX<_+#6#E3}p#;A2qSfJT?B81N44~&}5 z5MZCv0+{3&q-yPk5pZFYmsB=;09qdFs9C4i61s(@POVUt+|)LFGIJ-5zyB$Ha(u-mVh&eXi;|jIeoh zR+FfkcTB49bUs=VY>>udxjbG36um_#q9k-xYs^$y7Zd6-$kmrd663@l!U549;+*X4 zZv%O^511BOqPi!Cw0=TX6-u>=1)gYUpeMdKe{v0MPm`0Ap#T=Gv$% z)e{h0>@seuujjPwlhwF~^$`CdHUSd(K~B$Bm^hpcv|X0IBCQ$s@7n{&Dhv#aHOTpr z4_Xj*)#x*1kmd=7Dc8dSaG{A6ClfPJw*H|`T_R0-iSWmhl`U}%@pX@1tH%R8L%vtd2(#t1Ni9xQVM?!Ys1{26)7Ab>l&Dc}aT zqL?v=X{IrPm#bq=Wr;F%=unhrjOLD}Yj_v8tjsD!Yf@?%KGb+g-#W>yb0w!s*B2dnEl zQ7ta1JoP7pK69C!ZGsx9X=#ib!)g3w&SHVUFO^T>cGv>!##)}M5mde(b*#6>YE_jD zxnX`RE!qoifS=MfU?X|#8sm*XR00XCPQG!6JsQT{JVpgXe3J>w`2}b$UF#euKM?Ie zo;+uCTY0jSTWLB*(U8UYL?HVkF!7T5fmtWT%*?#g_2daU!v9tH|DUayk1P4!16<8s VfCxiZe$rzc5+bs~r62XZ|3B}t~ z8TJ0&-uHTc?|S}upS7O#taTQPHOx7m*!#2hzOHK@Le*4cu`$UoZ{50uEiWgfaqAYU z&aGQ0PIu9PPZpNls&3tK4Um_5rsZz1_4!^M0iu2@z(GlIv%$m4otW#uH(SFJ! zzKo)gBJMDMC6f+nmhbPv@N#LQ{1*HcM6&Ms=lEgzKaTTPdNPFiWQgZYh9_n7^sY77 z4Y9GY1Nz+P>89=5vRZ7hTLyZ5qFW7}gKU=AUxkna83R{R(*dmudOO#B;2IN|G~!17 z(!NB>uk-5{E{`B=qF;Yi$}kS-{Pndsi~lpeCW}3u$gaBa+FVDJd`Q}g-SlLCrms)o{c}VXUZ_gt(A09efaMZHkdl&OIhrTW`}(_t z=jDF)dU0V!&V-?FdVy-Didg1Q)hk}Id_q=m(n>vPnfQ2~Fb<@lOysbihu3*cE#-+h z=JwoYU3MZPvEGVDi9-n{BU!^2HXH{R8!BBcBwZX)B&Dwfl@_}q6@^|b{~T519EqY3 zh9mTp%My4ejv)Rgb8m*6)4HRn$~jDOYHAYCPj)jMeZH;-5*uU%1qF2+ls6qH^jMBO zs=GF0nDuykWW^lOy%vBYrWTHtWUt4e9VQ(hxm#XOQazCzvB?#o|A<@1GN$!jXkvtr z(~6q1GB#Skvu`cF;*pPpG~z|~JMgvC)w9cbI~gJ`z#Yp$$o7&T{EE7?6L{wHcVMs zxeRjNEKJ=c2%O8rulpgu)hR!RJ-*Zv*V*&b_{*JYrna^K46D8GZ9fMD=95+ynY}mQ z#WX8RBeGP4`ch*B%8FL^VhB+ynD7}@@|wLb!pQe)=Y3dvaT^1|A1!!fcptww-zX_p znts2z80xq+Sz#h10?JWTjJ!Z8(Jm*NHPm@9kEL34>TT)3O;a0Q#@T`UgW>E)FjY!# z1e;!+TGia+W^~mW`_D&2?D`GxstAOqtBa%}hKbmo_xaH#<4?8Bva#$>ob^;Ly+OP> zU3RU`MbrFthCVf;6yIxU3T#^aP!l^f;Vo#BI)_ zn_ws}<2HEpEpB59Z(-S2TQ-9xoy)89&2r`C(}=U86YxCYT3X4*ZtZC6?n$_5$R?)A z=Z7cxT|4X6w_lgOF8f~h-80bEG^4>c%Sm6eO(=3Z>tY&n?XBaKb!wm3sPPebMA7nr zQ`r&5+9QjCpFNFjU$7rAQe}*9`Qh)(y3O!vR=uT$#KCt_^2M@+atKfCj`DpZ32S@ex<~R&r6-E4VF96tdcaxZ}ar>Rbh6Jj$ z?vQp()RiAr==OQWVi_OJ=vyl{&fU*nG<$jUHWRT!;&s3^mESLn-{=sjvK%y$#Y{_I zs+pVTnR@SIJE6o#e#UuZSVvO(M@MZQn0BmO6=IjHwSEVF<|f z%PL|X@9w+f2{`1E#$9Nuum0XOk`tWV@kTfMZ1z?|udaR2hlEKXZP>WLU^n+W^|vcr zV8@J(D1O0=q+oDLB)LGC`9$f9IKi5(-Rw8|pCri`t8z$VRtqjw2Bv}uH|*8a`<7Sr zHlx!Q`{VgN1avHSI^=WsWi3pHv!0oVE_RZBU2zs6lv}zNg%xmTR>@J!zm&dC&XltG zgA|8=&3I{t$m5eh#yX|%^5lKl`5hsM6OBH)qL9-x*pp)9ju@PL$fB1(&{DV_qrrEf zw+Hd{nD#r^1D3ddEWr5B3a6XX?DK`A=zHOQucvzb3Qv_Lg~nY6+nnQ_u433(xF+!) z3G@f7$9}ykDLQl-p_^zxjkx&W0On%9<-1S1@3c{gC%g(poOhkMt4q;DQ#XjxyH>@_ z+w{cJ2hlT+5L#bSHcarb8d+yoSEN3K2~{@gh&lhK&$veMXntFoHR>tE(jpT6T1dU5 z@x;H5tY22(tyIQWK@a+(5XWZj2OWZjG~(LY+KekG^_*w%R(DkC>PUWs-ce=R;D4KD zaWsus9?_s(ZPR=FnCL?7;ydEVaYSBB`u0p6!nUA&u-R|Nb?hx;{n6$GR-^+@YYi9% zNnlc~^%j~}a66OAu+z|0_e(cT-xNxUax4=+&@m71`2?A;_Tflbf@u3Wys*+TW1Ib8 z=))}Ldot_B+^4XWRs!B6UnTbqMsQc zIP9EU3+{L(G;waaD#(+sJcph5YNemxxj#8GgnjS4{?5Rvjw+1qz-nUS9%Q=&th(fU zxNMQ`{e#AY&^OBtKkvez>&ZP{2|v1q7ZKC#o&8CRa{8Ra=~|-$+7*D8H|kRi$NJe) zGswr>LQ>rK)lCbXU|s_{Y2wkP8wp!(yrZy`ZAJRQo`O7o-+P2YsC3L#N?G zqk=lN6NB~E)y?T=fc;?5jY>h44Of?Pn{|l^+QbjD?<&Kw_2Xh!cAiULVY67*U#8hy zRXC+3k236{D>gF+h^;$kfl&KMQvH|l2{W=HzBBlrJuTnd9Q=@pV0==~*`X@8#Jazh z(`qkyQ_n4j(CL9+k;rQY`JC^z++>Jd{nXqiuTWyp{E?@^0=TMkE!L>$GsV}q()AU( zw6n)!0~qdjvM*}DxbLG)ZQYs`geIswO3BE{@efPYT2vd(RIE(`8%kfDy)FJ?=tXt* z24Vg--3ejsEqu6bNn5T6Ea!T|m42$sX5G(fbGo)%uy(}%F7}9YQ3~n%O);&}EA&uE znNN-I{Gh{2g;XFH>xukmbujI?VJRXwuzFxmBEOVMY^~&X-=GTLH!BCZED|bGGb^u7 z!{d{_+N9RObYc8yE-^&$1$S(*l0P|C;eVH51nQOrr*fVjWhc`>4W0^#JCilY!23e{D^MJ6YYk- z*281=^&z+Y_2pR*z5F9zmqxIT{=^e%UD&WfV2d1g40huMSM0Vx(ej2xjk&P4 z8@w}uu-g3BJQBX*cIS@dUF5+HRJ7K~P4Wx==(evHrN-_0+f$X+$DK473573E8>Z1F zDHEl$cosC@8{wb#dVdcJbIe7FQ{J>j9aOsLbc=66=_Y(io22zpQQ(19>@PzSE<=1@ z4yG}$P|IM2zQhei{}klWW&FW;=r$u?g`_@s$Ixef>XSvGJ`oUQUHIe#=A-!6qPj#L z8G3QDg)egGMn+5{K9`w)v25ioZu@>V8H{@*Q1sW@9XgD%-^hv`=_uFvK~+~R$yA6- z-WD>Kr17yd0B5Yki=L!(GK4#8)H>Gw9*T>&4E)K=Yq)I>v!+pP+|%O8DdOSTGAY;K z{M6T85$un+2rK&q)$9tc7K97>P6Ei$Bhu-py1(ssvG!J9f0qd;qjn)p;x(n(dG^dsTYxepbHlV#}ZE7K$Bvrb@->S4Y@JMEK(L~IYruQS>*2h}*CzoNv^ zr`fY#AyLre<1G~$tvhLQAw3cjG!21!@a(xH`npz{I6e-_yQNasq)V^W16UFG<2znw zr1o#W-y@yno+Qr5LhJ?=^3ip4NK-P3b0`#be0n*6H{BYkmb23x0-EYgVC_V^Gu`Rd zdegOoYgC<~`SeA0Bd3u?Bi@})xmak^vWH>|_*u6z4&bJ51@1fi=samXh<$YtX#7l( z1k(ADV^DSrA4Vs0;840fs46QoPIZemkX#^R$3g77K#9O3;nlaANY}+OrO|yi41$`j z0ab-CSzc;+33jW98CiGbDSNI=q(X;0bQ)1#W^URQL2A-yT_!3oc1&Gyjx7UsJaXbV zbcqJig5b-KVWcJa8O9`{b=LsGG7CmRuA1-X0Pc#TKmxHz|4p2BE+tna?lQ$H69=H# z6tm(VIJQXl>f>)LcW|Ga?Kh|sX-5x7qB|0Xp^t4Qzy8|8^Z(!13fArhN^Y#cz1HuQG2AeT%#=j;oB@ z(KYJ{*H;2wU9f>U>X-0b5^uE}0k}=w!UF!7LfCaHE7h(c)aNjDewfrEjrymRwY3H@ zD|ux)tCtUhyqizrp^(d<(CRQy=;LQHDyPszT>ed0~-G&%sX_9b98OEi+;8hCPyu zkEJ4_0&lU$s3Qayp8Xu7&yjY0WpnG3j{Hd~SBr8ge*so~n*E28^5&E|?-PB8+oh$Y zdEJz6B0NuaU>sngwX|QeJW1_6OKUhB6t)6fA#?~wWbc-~(*WN>`JH9ASr_evPaNw% z`D1NKMF?7Tq-k)+ZUSj$Pn=9ejbcBxX-patSM#k8e}3aeCghlBMaeSy97i<%v^1j> zKw~}^)AI~sCCjnShc3r2bafSL=R8@{UB6ku^A(AqB_)dS#*LZY$E4{VE6FD304lH= zFMY9K;)ex?7i%i_siX<2zrr8}&FImwf=jbpx4(@r7pBpSQ}Dc0%DF$kTJC$rdUbu$ zqNd2j4tA!_)%ZPdjZvd#Yb<;z+;)Pf-R#)%_Fk!cxy^i&+B7429+0=deBR= z(%q8_tDkEl3p9`00x?_giQi#y)IH`FpS6!Yqymub{vGHbMzT*FuHGXm@1N@DKgU$9 z`V!f#9lq@*uAh8=BLX1ph0|hSYr!bM#H`Del>4}_%9M?mxuf)&SXrblhs$S^BJ3EcL1(v)=ZTz#` zF(iFsSD8DfBp=qt3UOeDEe4ND_7}Uu?PL}m*P}oUyQ^KEt#rl#`B`K|ui_pE_5`$L zJ)ORz0v5VDpJ*Xc#UBWD7@1o2SWW-nPWDAJmRk&UIzd%%4eyHun<<;Xb&L0f>%Mqm z8~%GjM`1jz1t^rI9nWF(ib05qIAk2A8{X1&_IS+xiH4>oZ3G#IT+2D4yaVm!Knh>z zf%E)8$oZRyl%An889ueCUrr$=9+!8aobRksjLtG0k6kr>bvEQnbknl;1?%U2H&yt} z{z?662L8>>GN&(YAIhvLt2hqnlXo1ATa_K>EM13Vf~0+)elPfJEb}`4?v6}fnzbN& z*Y>CdS701@tl7{I66UGSf)`t#)?44+F7R;2RMf^Tl^BDakE4s3HWNO*=!n$dpAT;T z5no8L7?*Q?1rKB8#5+~#+WuY;u;L?vmc2exN152}1A=?YM;z`u2c;%I6(6%#z~C@g z^K3=ai`{87KNsCN0gJER{x#UJkap!2k(l@Z!G%!%ubsA?uF2onur}t}C;q+*phC7N zR)dK5lcF`Eo?c%o&tr8~+lBrF_V9 z#X#F~F}tg@|HKme`Ta<`S+~BQZ{WLyK7$yD8%uke(PyiAJ;%DImuxU67Li>6;E10e zZ?jT*?lI)Y09=u}vofh^p`%budoWI+@AYA3!IRL2D21oW+m!ZA?8Rboa`4^vcH%NK zIygx{`og@`a8Hf-H?OAJlo-C^*kLVKnO=IhGJsFaapzG^vKPS9IQOyF4(CXR!==#x zn{_@!i47r@$ZXnQ$oI)q%G?)UxbqNWZ|=?R2#`6FnVw55pUnGmE-(+T+UwzBWLK8XRd%(f1x4VvI%4I?_e| zP^|kIX>@nZq&cVwkA!dOEk*n?qC1V(>5t$}U4vl^`&E)L@w^MWWS?lj!-qPRYMi83 zYvV7@UhJ&G0otoWS=-x*IFp-dIaV)vJ59(*Dm)1@u_7g6dFoo#%`2iWYxBYC!@+#k z@}^jRyBW*+v~!_Zq1!XyUlvSynG3cKVrBsHKS>z|I|?~l!PM(c^Dt5|anhj8XNs-= zTkDbkl_Dr&^Zq!H$<|6g4d(dXocIsnT^w8_|(JV9gTM;q=yF{nu5pRX1pq?sd@fGneY>E3bSahrU8#njlW zW&WS;;_keyE%D~R<^bpsg6RJnKPCtO|45(iIqczr)?W}|mg?jX#re*?g4VBsp|5VV ztC>ub=_a->3RJT;Jm7_#r>GV(>OjYX-$jb_&1;0OdjPX0pqv+J-!m87@bK&Tai%vI(g$l zTqR+fgK#-}fV?s>zkTpH62J=o!%cTz7B1M|JfGtfC!Tf7P;W18*mRK^vFxgJo}6mN z<=2TYCD24;L;Ks_Ve%yex2qgPL9>$_^%8WF2{6t)Idkb>xA>TVJ5hkUoc>?#Vny7N z-mrf6BNY91@XJR}B=``deWS(}7mvTe<@nrYzM_XaM43SS@^5YCTd11NO&~byOJB<@4mApEJ3qYV%rUFbI%)ecUAPx*hzdP&V z5u^5IJjB+~3o! zQ)lJZQR?~Lg}Fvo`a!HT;m>6U55p$Xd{S3|`t;skTN6!QZ4%HKf8VauPvy&BzOX0* z16D&H@PEnU$5^Udc=$fwFaVi}_q8sXGugNz)wsf+OTdz_xb*Rm1#QvFAbBz1>9hrM zbU7S}NX*>*`Z*fgf9!NV6a-(Sp06q|I{xvdXyM=xuL}wtBZQxJg>Nhj_+t|-lH@~B zbTR)LAv1QAg&?Oj$3yF83tww$CFfl+^p!U?Li~^fZDk0zZ;8D7e&Wz-Zuaxw(FAP+ z`EClgM-&bnxXxehF%BcN4V96Fl*Hq#vM+TJu~0wh|8Rb0OG3TcAhLNfIZ!%;UJ2B> z?3?oqW0zqbevjaYx42}9Q1(pWg_$Y)z`C57dR2!2CKf|MEuQpb{3B)9$|G)|x(ajVp$3%mk&Nw>-;c zCr5xV{b|t^9mii-sO#&eKTT!E4=YAZME%>x*S;%*&Ih^aG zIH;zPcM5{o@PfX%L^<#kj?D3CX#j`+C~7+&iS~(iHJMQj*J{NJ`uqO(kiZ)I>;h_g zbp*EUfCEUhA=d5Ef}j|tuFZ@&xy}`({uOf#{qJ~~HqN3ENS8&WIiP>Tl?>hRbM=|9Nz_=)Iz_s8PM0*uS7_{6Rw6m#w6-(w#XfCL|@qMpyh70XLx`ov!zTHpuz zZ$RkZ=AvU+k;p`n(RUrkIB4JX6bB$e0{x+)iwob?!G{(ti*PvZx=Gu3nZ9CSx@!uT z|7x1k1NIk{9UmFuusllq%u)Zgu2Z!F9m?l%ntN;U_XupQwX~Uq0BeY!^{7h8xIG@Y zFGl6m&)wB|(AyPBzP|NosOElNrl12&wf($7I{VsaJ{C3!mqcE|+H9kM%g)RT{dX1e zPOL`te{V`$IQYlan;ssM5NW*O%^sffc7FyG_B+KKbkSL8Hb)x>0?fd7Kc4ZQB(Q30 z1V5mFOCNS|`s`z!e-FT6!v@bcd#|q-W`ZrShCKr#4vo5_D4e@UEtnL!Tzm3S6M6pj zwgyNjT^4>fgaO6YYu5$Sb}p-IflNbRj`Ep}8c-muBZ(y1Xb}zPX_wtOedixSXhDl8 zVDu6I>^Wbbt&XYso@+J$P}XX7a7p*BB`5&*j!#Qw)YFY3DbhId2EzBD?SBR2IQf!N zd5I6~*86l>fsYY^b$3_;0^14-l~Wla-~cQu2I{KNN-Ch}3%!XUI08yOMX#5Rw3=vU z8h@t``H|kFJ5^VPaWIf{%AKl*c@6JHOw1jRcjr}s`m^e(@UG7G^j*|lTv`Sppg5E! zW{|)F7YK6S>04a-Qxq0DW;ADt#MJ!s^*?rba;K!bU{=38H`4xo9^og8{8G(+%^HXJ z?e|lGOA?PCX3byixG~#r_N)USs1V@kyH^;j;Y=f%HeOib{-syJJFz6y9*~t_$ENlJ7Sp z4?+@pta>;klS;%df8fx32l@xbwKU@%wO9h|igPb>VF(d_lX^&#j?&@x@Pxlh*FZI< zj5|N47&T)3cHFbx$TR-(x4dKmMPll)gT2mP%71$eC`}4a%~UTT_D+0tuL_`n;g*2u z3cS8!TH^dGpk_Lb|B0saz(K)Il$Ypx!Ky!r!X4E{IQ;KNEf;|D$`YIKyVSO?YLk@+ zWBvHCeF7wP`Jp9rOad^+HuZqc0MQQudMMlfoDa8$DfxX*;W=959~;B`@Q5K^dQ$hy*InTsi#%WzWVsWmd-5z(^gy zNc;aWlHD+ihQoW2%&4psVLM(R4yFogjLGh5x({)2kQT@Oeuug4l{tsl*`;y4>XW25WD&b*JN*5QnM67h!4baB!66>lCEG-uAOJS z*O7LFOxP(lW$MJ{dUyXuIY7A@^Nno3z)(55j6pL{05x^aH%@Yzzodj9-eeT21I!Fb z*Af@>-i4HNfvFdL{$p|Z)sh@6OL>xZ&@k{S<0AB%K8vFcnl=Z%5G?F%;y9E5FdIy# z!oR_FTLI|!AwDZ#=}&$RoH|Mvq*~#>mo-KMk)|cksRypvRZpxm}9~WOUg9x1hW&Sjs6=*E-98+a!Fk8zSQuZmn}MAF~u8h zQM+i)&rhtxtI!T5*rQ7*1mV!k_*$)hL+U%>s7G`07nq$ zmnzUy5t6=aabDMWdVRY1E8YU+E-R4m-~D}3oU8q<4x&iq`DX3;F-qp)87V>7IeUZK z8!Tndf$_W41C9V$BnFTnXk#NOfwnhf!`k^mIOn?_lNp@9UCnHiW#j1m?xwufL`m@4 zv0?|vr742Uz}vtXClNon%49*%`eT&HfgEOVJFoY-^V5Y8TB*trHKq>q5~cr$cmyz$ zlX9T?rQxDL*|?|af(X=_FBdp#NtGtROqpQD*Qhel0Q{~^4q3E4b=i<2HuS1U67c+K z`4r%76X;0=zWg1nG)lfV@WD^VC7>H{mMJv%vrH){+WZnH^S8;Y(ic8^9tjk(*oj*Y znij%Xi*jhhzMTNF(X()3PH>{x%-@H}^WDi#!2svaTmZ2r2O#WQGDLF43FIY{>AVcg zyMX{3|95!JXVlO}Lg2!07Y)fuAkZp5fHtu)ZKsOLzC7|DD%nEf=f7zW9)Li+dT32t zh#h$IB8z z(z&AAF1DI%B$%t#*Uz7dhl1IWlXrdk_kJ>40TS0&fo)c(_bDZ`?y|{D57gbF-fqLG z_Ig0UC80?}0aDFd3POswe*VgD$^mJ4O~m<^$tKwzabMS0$yp7F zv;?2Q(szFC71-Nb6ew`RRm9HBkY}x<36P%r01~c{t)~w6ANqkA3VTA>7RH4lrM3;XIk^l~S+%@{tOkW=%nYzFig6ZUAbn+CE-@}0o^{-e8vQur z?tWyd7vaGCk2jBvLVjS2H?3RK>*)nO?E&OSYr#;w&I3BN;Ledd!qm-!mb0f zR2Bfl=Y232N!~vdJSp+Nk3V}hp`$uABcY=vwJeCH4BBoE-*fcEEDQ54q*0n3!-FKS z0GkX1Ot~jO*{S*zklWV3(njYQULSmb0|F|vtPK1gDV0E}r^`Zes0m!M`%h>;fMnM>1*IB!}FUbk=q+OKnn0QG<;Onpm!V*${ zyVUuBz0Y$%+sO*WX9UEsAMJ`x)G;)CfRofp%KGLz$2@0u8^~@_U;2g%j)H6E z4PVK=H=ZY*V7-+HS-$miFX|@3dUx~X18_?-Hs8PM+xSOkqtHjcck+q9dnr;)2VfJE zY;f;Xoda#U_wiKX#lPU9*N3^!Z!}j2)2shIS5_Deh!DvS;Af8HXnET(?)zFUdP19v zy>srbqYg{vquvIeeDw|;LPpqvi-Coky|>;#8N>9V#*z~8bB*PCsA^YflMs~^<{WzN z{e^d2^jH;Y>EiRG3_6?tMe0k)7rlnS6-$sE;2iQU)E;JS{%Ah|Ks##wL0?AvYF%|5 z=;+P|x(voH&khm)0YQn70>I)d>c0hm9_-(qQ2}Sc%iiFaV9<>La_s$bR8*`Me0EN} z_;x0#tLqMt@2O(a&b#u%gh5hCEF(dlG3HHHlUb(`_b>QkD-BRS!ddX|f(&_?}24w{8J zBbkj05jHW4Z?(8CpCyg*6X=@U1Uw#=Qu27qXa>%2otcpYp6qIBJMaM|PXBKMKX?ny z2HHU)qePE&6c6L!d9DzCa=S})+yucpz^9<#QzBvxLaoheT zU6|?9fZRvHx1`hhd? z;P$ZWhIOZ3J1D*EVbEQp#x*f4+gUb9@!pEw>>^G`rP9SxOMcs&3CaC09rh-Opm zFahr?ISi{P$glQvlkwwtQ^sumt!mzj(q8Ogc3-A_4!tfca_lD`b%Xi)58s;LQaC;f zx!%>qwD)yvX~#T;S!sXvG-Z84-;CE($-dZa$t+MxmISa8yDI_Uw`qEJD`i#6#Jtbd zCdzsQJf0+Q*N==6euM*rJP>b{?nEe?r-m0DUr#etnQP{-owi$Z7rmyz0V{h+-Mf>a z?P3k|S$8A@?E}PAdPp%z7N8IO_NjGKUTQSc(y*G zy12-wSXvWf`+0iqWHe;=5x)+5XhIvL!MT{rJDAF+&fVr=eYrCb$e3y{b+&%d{_FB) zV9rIOBfTGo3YjKj8wh3}vQJFO1SCoj(j1GHpJTIB?(Ep(&~#(67T)^74Si$d(Fh&e z=N6(!(nAh-(PQ?_?&qH+hF1zCAAkd&(crQv0|J3Iaw4#p|BOPQz}jKD#KX|ejF*gH z_eTh{{CqWqRdv6b2e>i`g;-pv(V$BNzrBACZ zz9|py-_KHm)T4|83mmmxy3Dd{dvb8FWEKu)aH4yb23i#m)tY+GYxQ*cBSduppksEYs1a1<>mn`++qYR9=;)tD@)GRgSj- z1*kIzA==)=>-O9;?lye95c^_bK*RFN(=snwOh3uGNh@M{Sbyjjoaxj+QA8V8%#^eO zSq2c!)u*QfeO5&FbUXp^)y_*~j8JJTefM!mC z4S}kBfH4oflRuFC+uT&h;*o@HDXW*CdnfpJ{f(KQpoxw(wG;ztKF)3|2IQ64whao- zTqLaie$s~kITrl6xz;*isETdN8IWjf(3}Ce1sF3yH6ZY9{*_ZpK9Gvk%r~*(3RsSP zuXv(k0jV}eZ6zs6k0~B(=T|-29SJek`||A@#KmXevA(4zteTqxRPQ6pNa^7 zZHnfM_v@L6R{?~PA@F;0*=keQrejUu3hKAZfLX4uGt$|i+!`W~sh!x4JcBl!poOW^+0URbT&?&7x zSV#UxM`Vr%yWS@XD%m%xOW~MEQ6dZ~Z{y+F zJ6Aye*K$I@*$?>!)8(s4Ztrb`57^&2-oW*}{joaePY+YCX+G#NWkwR3JcU$BU0?8O)A}d0%M{MW@uD8#Ip?XHJpMdsD zZ~lu?;Vv8cxjwcYLL2oEC09s%8W!b(2og6jWSQ_oH-YgwH?3kyp} z6VR^`K*VWK`uFw}fB^38Ac9PV&+V%H74)h8GwgCPPSh#fVS@)|V`1|D%!zwB@{iW6 z189JP5*b5cDm_qxsUJ_cCmoH}bCM1*OG@e~Sz#nLNaGZMn2wvOkBpK4QJbs|hWZqH zGd~*6qVeL{9H)$??T~PP0d$rg&XvW9M&v`>=!l+B-G?Psj@0CtK`@cGOtzE6+4ke_ z)~=FKJgp<58y*~w&S%)?^kvxObah=W*EqqAfLE7PoN=$$WW^PdNlrAOgdg*Pf}!=) zmwetU*_x^@j)mh!b5yJP3=1KCaR;EQKvgdgx~P9fafMlns-OE`~3+#2ockHvXomKS`_$plhHXFBj;`Rd;czhIGWE1eW^Q`5iqqgEIKayWq?#62b;m zQ=lard4Lc>H~*8F;K);lmHc-H5kR>LJj{>0pyh3UNQJQd_<=*5fXbZj>dFO2RHRqd z-iIPrIk%O3p@s?j3+7dUp?BYLJ!hI?VKA?`=kwtW;+0&ooCaOQ}!a0u1Yr z8I+~$k_RDree0|_LrPsW&MMsz!9tAW$ymKpazW{g3MFpzi@B6naM=lq(CCIEeNNH+ z`sle-+^pjHV9TqyyL@Ni?Q;3;>V-~Q72RRkj_*IGGmEEaqV&Z$jOSv#Ck*px(k?vWTsxX z*3NnF`?DD81ZK`t6LTDa4Rd~wFIWlTokNHPd*KSi7?n0;M^b(_2^F8&OFz#CvzdHi z8;}m!NZt-vTh@%zBFXOnE4ElBh+fi!)WT+Rv>}>+K8|a#ociQhhdv)httIF#a(7G2 zyx<+n)SedS*{4Odqi&nV4V@3E_{>2fV$cz7F&P#LWO_zEL*YfaWD*N)ekmv$c|;ju zz@0h+r$ww7bc?xSY2;bPar|MYMYn==3Ima4#IkcmFGhZWcD8atjK&z09=ORN6?_&J z{SA~rMzzcjOZyj*u8!&1Ie(qXq@WOpg=r0qASeNM-bL1JHs;CV(7=V$W&9?wzy?`jTkPSc>og+A36cNOZGL?yCKTX1I?{3>)d{xz*`tw23?72Ykgl=Ytt zvd-;U_qkcFn!h$*LA#2$yjW|31AlA^eUZ)NxX;j=-lf3Wl_F`K-z978^;vHr7sBGJ zMP1DTdgStyOM9_9THAuh92H>2s@@p@k`#AV<kaTc#M{xAb0AhiH~F)_MBI)i%U4GAxzXPlx!v zmM&9il}V|CBJe2@}(x?=}Tya zXUH$1s*qPV>629;_*NB7mN0r&_)1W1Yymn7qw*+y*GIn(TEjVuYvJqT0ToA&3d zrZPRmbY%Pb?k&8r?4$aQs zhF~LwSb3!;7~KZw7_rhF{X>nC{c6hkLyq3Jc~Y--pOCA{BOrtalkLU(Bw_$iiDcf- zO=+~xZVmg{Ei~rH#e%4H&FL-SB&tyQ;r&Kwxz&1ptw6Cugd&CMQ48xTR6B3smPdC~ z4nUS#eX~Q-6wSM|OGuh^Zd3pyQ(6W2#L5=Sm~(-~t`uV9u{BKYcMGOm7rIdJ)cTp4 z^R(RSI5f-pR!6Y?*>-Vg$SKPfMgeB|FnPcXk1f?>6)Xq-iKZ3Ym2x#5E0EpyZRIMF zFAsZwBP9a`d}IwhXvFj2@RU29NZ3wL_Vue|L8S798$cf!ZvC<}LYX@Gzx#9iHa0ej z+;->S^-dNpxLpkF1}M4e(9h-MKpRzunWMEGAfi#>7e|4PD+EBNwx`6q&MUX0tB(Z-U}G^<k;S*?oxh$x**)Cgae`8TgmqP)cGz6 z*d@0P*J_^%`Kp)Hu;(?~0aV9Q9d7>j+HV5s(!q9j7)ZJ?^_?y03HH^?^8_>mqt?!! z6F+Y73Razyn>OjLdhgMzkl^^V+%exgO$BSD*51$)=w3lgp;j3HxjXE}ZENzGMVdoU zSO!hJB5p`z|v+1WLKESwkkVZ+$rV1*9ZHUet*N3g)P zSP*u6no*@ujX`KK$DP9VciS*|zJ@ht8^1Z8{$X8)Fn)W2W0Vc}X`rS)T z>Gb2|V!bFgN+BNL#$pK{zUV5s8OfkQBrr3`=8zRSo9RB74Z-F-QS}orYLXFw-*7eT7-}csj6uVIv+vOGM?Xuie4@jc_8JvUwTABL^}@7illU9u!%QQlt9ypMb~GAUvFmnD|taOm@al_ zUKNW=?AvAI0A0XJVnlM!Bt{KCl#lxsx5>KJAsyqLtVMREChD+`^!ZXD0;&WC;T8&H zIgmqA6}EHp?#rQS?r>?QhWptv!I2%$2W>C;AM~s0md8WNhmlgL-X0$gViNOyHlb=wPBRzJ*z~Nrr=;K4goI zE14_=m4MZ4ZDH$;?RvofI*CNS_0=F{%n)e~aN+V(QE6txEUn5GoneF~!T?*8-%1^K z>%k+z$G!0F=Z3GS%mH?<;K{39fKZg1DxAGHE&EDqfI@)5zdMovquYD7!NtT^X0gg) zu`AxzSfLilBQtru76;A4=VrWYHrPtWavM{)asU8>PS6FSNA$UzTWIR3k`8}7~Kf{-|6|4_Uo`=1S9RuAa zTT85kt{GBo<$e&GS7~Ro&+7R!F}#r9r&%+k>7iW5Axqv;IGkpvRwd?og)N|AHcLTZ zFa_3Abg&NvD`gEg#W6w75|}W$E~WuVrL$@8i@2KE^e67z*ypq}V zH>5o7w2K$lKU^`(o5fyuS`~YSsoeR5dnNa+vsejC?gC%sk`#2KCKN2Rg_Dxa%kWKrF_K%27 z$W^^ z{z9w$`IcN(ddYNk5=5wWZIu{vWm2Pv7vk?MC%61aN^qGu-!B{RsX0oHA^;rG@p#gd z7^5>BU2gCj9)U)QMj_C^gQa11BHS}nJ>MBITpT*|L)xM{F**4(Epji%sh^=&LLu3g zGw>xtdx(_HxPP9-`wLgN5n^IOv)thG+ESb<(K=2m8nNDE$U}&~rm85&8T*UgK&n9W zPSfF_)Qf0J(IVfQYZJ@N-*mb@woM`f>AY+B>{dyc4ZGbb7uMRiq?#*aw25#vQ(X8Z z-nXs+KW;*a{vX=<(=0JGx7udi8?jrz8Vk1(Dt79)a}fcGG2O%f&pITgG(ywTbfZ4G zQg7?B7!b-~yQMpgz`Fz{TjI{RG?NtF60W#Al*7@09=nWq)sP^mKQm?YB0zw1pMOc4 zRg%|I#*gx*U-rhG(YuzCfzD3)TZCHyhULP+dpWnaaKg@a_>lSplQB$vj~JKWo%J>a zW4|vpb~(daHxvSL*sKGrOnFI6x0%gHUiYhAaZLGFU++vht2Xt2w3o#O1HGVD@FMjb zU73QHB)pan(}DJ`-q=8tl?@!<)7waP7T(}u11&C=j|cvfPrQoIG6A;)-yD;aL8%7g z*tA`CXd9!gR}gY1rCd|MOTxkYGS zXufD(h})1`1rYOqRu+W_goTv;0|q(U>@Ng5r|HHp!7lIiuR%3-QRE1IU)*WzvMaG> zqyAQ+X|$3gfG=CN`}lsvq-UJVyMm7C$*}ldg=$`v+CW?yb`f@F|z>+=$03^FgY(+eK}X5%?Tx5@c*%g9P&(g2DT?0!b({YxI`p{_-C&GgMB% zdU<8wq#|U49neJ>X#9RwpIciQpq;UCdQMC79gj?#oW&a$xDZ;W;%*`Q@-D~=Cm3bfc*uGr|BIgm zIQZBW?<&;aou&R=NcGFikDfC?6Btd)uFu`5f%47CdP&R58gwr=Vu-TR<`)UetWGG?)X@2zQRCvrEVH@EJi;~=TiAA75~2F0veTB+p%CV zp+yszO=k?}RU;D#W!{=kLx(3a)x{^ zcFDB0LJFs@VsQfhp+c)PoYhDz1OlW~Wd1B+^MDgxhYpW&oQ9m3HhkGY@K0$r7AVcW zeNgF*e8T8gx1WXNFeF^H$A6?2|0vhvLWY9a6-3HN*JM_}thP}@fhl}~(dTeiD4$;n zfM;CQzfy2OafMeQdG zJkwB$VG!5--6()SmE{j}Kr0oT0}fBqimq7w%_jBCC&{TR8Yky|B>WBumVin%pvaVu zqB8rF)e1Q`)q0t_3$Hty$7E`{m* zDI=6u8#EQ`L=6eoKkm8WdV6FwiSyO?Dnf&Y3$fyYfI?A)_AuGSxqqHK?4H%*?b$0- zjoA#ugk?yJUT5$i#D8GkquMW9Ectw3m`U5HU!6{xrv3qjPGuw21*(jOegAgpEc#mr zY0KSIy<1qZd96hXxFiy>jrU&48Ay42{@pn+;1=Og_}KNGy8vRtL`=hSz4siD+#0SEpi5bJ;bbh5x;uyt? z0x6#r&*Dh=#Nam1U)MKrZ*dyEmz(QgUo#&-9|J02nSB)1wBZ^hM&-uFa1XMtm3&WdD8pvDQ5 z!Saz_6T!a#oQ-^V=N4}HdG;?_LmEi{J}36cW9O%|3uFVYL0=5nj}2N>%^Vuu_K0K~ zpyQFwcmp4?k3eMfh5A#@Qp+U8z@SR&?sp1yH9TQlkpyAcg!KeZ4Yq8iaP~ARFxh z+;Vz%Ymxf1V<|)|eJ5gZ)abnTgSTQh#4Q&XP;?)_r=sXtShn|M0!%9MnX;zPfVIeC z04ZcHD>dc+$qnaFu^EPfPV@gn+lJ4e8emZA$1&H=bZ^ZfGa03UnEJ>AOeTKxXn|1g;@~YvyA6T0ys?0pt4!Z?sxamV@0;+uhBQ5#Tj1#1A8>HHyTon z$XIy-v{ray^H=unVlZ}Bd&dKai7G;L>66@yPLjB|_XL~rS?3THN}$27wplekXAUf8 zj^DBEP|Rg}!vBOQqdNY$_}hy3*90OzFm^=YW$EX6g-X!PVCN3s}$E}D{UH=nbnggb^Ch@ZFX^0X1=wmQiukoB!?V{Bidrg#5u^XcEt(>S)-#UQby zWD^C~Al?qElR{o)Jw%yW}dm}5ioXH%h>=Fj#Ql62^r`a64-mp@ILRJ^8?lmBB%XCis zxit-yQbz%>l^N&(92nh(Pap3s_ZDbmGmKWQABWyYlY|5bmyybPAI1y;$mdWW169^O-^6-O#n_7- z$VkdP35XqCf3i5~>7lGK1roOZ^1RHN03U~`nq(j8`XBYj@#4zB*-d^RxjJF*0TcpF z=nw~ZMbqa;lBIf)F}X%UKMlxXX;*vTwmg|2Vh4mpkqRJIylb*7L<W8x+|2ePAHS+kVxt9<(;#YEWaF94PjLM#VBlzkD)X|`oLkmKsh3v5_qMg z4d%#{ycnHDw<~JJ+9qtnLm}mR@V&J%)~mYLTXVpJ%--MY))Y4R8k|KE`T)t_;g8lhiS!d_@!qfL4RoMn)d;m6huzYx1^BbBwe5B}#|2@s5z0#qLM-V?B!~P>~j^Bm9gTGbj*6@Ez!~Wd&N8 zDi>qS`_K?h*E5#vD~n&90&FRa(BDBzwXoJD=9GDCIgT(Hq|4D0BkFEfmTyWX)}6f zXPp+t@tf)-EplLjtzW(aIG8IQIxR#jQW8E6YA}TivB#3}W9Rm}_RzpH*h#jDZD%}(?Cv@bs z$;L%0_j2e4scse}NtTh$d}zDpzw1bm?VhLWyY*AI#2UfV!6|1~KD5n;X8X47hPv^(>|)d1O-aK7E+)9F4naKyIm$c_80CuK zjE&Woc2iDvorhGUC}s}2Rd@;&z6S+I(b}C!YEcT=@y3@-Bi#tAROhxkzDYrj&r|2& z6!Kxf!v@5|OJ*o;?thMvCQGVA*!_CauO|m0jk70jkOejze2p2iZXl5L=sC}pbVCqT zkN9OZue)TJfHS^8i!WtcQlgVxH|J%^$>9lH4n+DdL&Qn73rZ;@l$GH!JQXElH)^hi8op&{D&0ArZMuBRdf&$*@b_C&X zxGB{sQw7Xdaew;)>I?>;wJXwPiW9{+AgCRR?T;;c8_C#cH9C?kIrxA*x0(IA??r_9 zg&>whs&7(_?+lJts&_)*)T9zK>|9=Ar{~&qqw7@uI&N4{Ur3@hc5KMWNvEAA*Ogm4 zUTAjWTbK$Abno%SOtkg~Xm2yX^0>lGrS=T6e~#aplRde#5bh%vy@z0>^?L3*6{%@* zyJtqy-$A$!q#E<*{XpXhHeh#8^EnJ@ugSVTX3m(y((~qOuD5N|;W%&q?Cs7{fUKaC z4J;)=e%UTiMaBeh>P$F5a=4~@hEmnMW0=z?YkBQAu}Mdaz0a79&BBHe=bKSRIZH@* zZ%>ob5{rX_RSM+}bGHNW6dkLD&A1iiT2S)xoRDB%W?~`W^_8G`s-yvYCnGiZVLcuZ zLR@7#6U-1wef%f(f`Y;w7_>*AU<#t|=Ny*b>35 zNKH)sX14z4adzv7p));q!#1RT4R%JD^dmIuIt9mDWerj?I;WtLyxCsPZF#UjmrdI% z^?|OorW?fh9HK^4i=SqTn5nL9-N8waHZX6dg2`P>91PA;bghp69Jp&lu9Jx|zOg)1 zr(0+ACU`G{GBcZo(7qyKZb6a#2hR95Tx{Rt5Ncy&Bp5wwlEBlWRq?V@CNTK@>fiLP zkE2(+D@f7JRA<(sV8s7Ej9Hjj)ZQBtZJ1p!l>lvdn1vEW7E_t6%W(LhzU0i5njUec z)!-OA5Ff|W&z(iFIMQBQrnq2ZG7T-f0iI4L4!cs@!5)#$2nSo1CN1Y z{CU&NIk1vFu9*eYF3!1T&&kE7&q|Q?ls7g|_{C0G+Xlf^AHbW^^4V1DBFQS4oWdUd zen+1a85H9xr%RK^!6qj@o$N^dqK`1nPs@W^p@k^VMK8Yffv0V6Zz*mkD;9>Xf$$9$ z)%40~9|2*(8{VG_$H=ogYGf>K*8|?@m&piI%E7oLC_wUnXJE%;6-i330`5#FVP)rd zi?~#7qi%Pc|5}mE{?cF0d#Lm5L)@7#jsnO{I=w zQ&#SylMx#SlxP+L!ag>LTTYevA(ELjZf$t8E&JoEV8-^>&bP?A>+PBXyWk zwxpL^DiDkBJQ=YE-yBP|riW!ktka`?rp|MMy6?J^kA>P3yf~nt7soF6Qq|4}8JjzD!EKjKgAf242IEFXCf!76 z6~xPI71T@SL}OQc#8MnK@g|yr5nM(?I$GRY9x+Bb0c!ny89^kBpoQeU&tv07rLS+cMcc!mG zqU@v$b#G=m&eUN!G`!r!MAg!$sq&}y{dg)qb6#gszchCuuLLnV&|Kk;v!!Rt7?Bu57>ax0rp{x1Q`mbKgK;c@O_J+E6ND_7nC5- zl&X{aP7(G5ZFgq_XCgBV%;^}wQGZqlepJPF>hp(8x7PIZ{$Sz#(RxKvIt{!P?cAlaJ^|4AMB&fz9n z+xqat@IwIa&4;6eTj| z-dzk6e|sCogn?zprHJQ-^sgPeNb7C{*owuqkG#*>>@F4oS205QdEcx_ijoSGBsbv~ zOvX=gkkcdePiFTE5tU9H;-y=**B*L8%_Q$abm&mR=9B}2iSmuoKzLt5z1vLDE1Cd`A<(4<%7I+L~2(=f!P?gb#hRMIvmYVT_)NX z2Z-NNCSKY-9pN{Z!IbIf1Y1T>{p5E!zNWOZ^2AJQnAusBL3vmdO@m@b5c=Mj!edV~ z{mNmUOL9PsLE^}Zm(h4mk$28o_i*#2&nSNLubr)FDBA18s z8KV6SbM5MdrvXR5^?==EY8njo`j?aj(vYG}lYZ9r`hd9`}UnNiww6h7k8Y3`k7m*?1c5EBm0f6~Nc5`%VlSL@E-N`arxsPPNKXDv zc>-vz)CsD5MSlNDr(FaAID2~a>`LI&)_v%Ls~#Jn66wRdhH+Z6)QltEd|tkGmQ4I4 zG6x&;BR^xevuhRyKI zMszfW6u|BEAJ4K>jP1^vB2!&@`yM&8t0*(&@QTibtQs>&!asDhOY}T!uslUyqBJgX zYwkJOEY*=(c5)CAZQ_TwY#B_#CjRw#!2F*_YpX^3smUCe4EK}B#izrzXw<2Tzftd& znUpeI7J-K?m;c2f+Y1}Q>ZWa+Y*&6f z2*nt=0C^xB_vR;Ex{TluZ3%@kO{o`7R8fb#_dK{xB0u2jMo_v>ztb`WyBjp`U{i(A z7;O0YKP|8mQFF_I1L9EDQs!q}G8Y9T;h$pHMn~xKhKec&?k_o$c&}CpZwX;Jt5JtiJs-V7tkSL@1<0&r=mS3V8$>S zGc4pWNEY~XJx9$O@#U^XXAGhUUiia0c#HK2r8R23;#tNldl*!4mVv-p;UhEcuV7tP zcO^pjx|@L21YA3|WID5Vadq=cXg^NhW(^ zi#w6fn5M@OQefk4li`P-Sy*|dp1NbZ?8w}_NPZvv3O*u_%s6BADz%upv(3Wub4nIO z-{c!JrC=BR>WVMI{eO(fKyvvTU;gdWjjS)NdC^r*43p%aR=Ja23Sl}4low{b>nc2& zWtkPJgDsD6vVsPN>mLb6zSc06{B>=hy8Z)_m>=<<_h{+@ROT5uGHGwS{&~>97F#xg zUY*-uvrG=qox}o=5tYi$XAiB$liPa%OBvOJ5Dq7oHG%xscPc(RiWrUw7MdlChE{nEs{*wb#l8P5Z(Ft2f> z>{R}Ha27vlut6pXPU$(00a>g1v+%WD*84$epXP2+YXtkQ(j5JNdy9W9+^_$e70`-AazdQCRKq$b;!Sm&J@*}|9od`&n{nzdA#a_AigM))v z&Bkys!8YkSB~(6Ekb{;>)t^}b;Vx7H!AO0(D(odO*VF zTZOZJ!G0yqzPgpfH)0M3Elyxm7=rbF+pjzh7rxv8H-9V&5$_OyWLOJLp8?4JEPSXD zk<~I>9!-)e)zi&O4s6~5tM->1OPC6S{jiO*^CYXj-^2CTphD77lmxhH9DR_df+kn!I& zMqSX>n_M|26^-bG8#ZKN5a1mb)(t zV35agGBEemoJtON`pT7-b`IEusHqJ)49PVan3dWF&gW|$^@WYHkFTo4|GBgP(Vf= z|7&Ien;Z=GgNA_k!S)8QT)#;;K=}|tfWiC@ z9N$Qi(hJ_#=LOJ^{kF5N>vql^l}UpTR!kcu5LRO!Z`OAP=*sbZoS2{NHp@&3%vw$PQpvPKm*U`?d(gh$Y1niRrF!JwhZEa24ZI5~oRxpr=LQ%<_ zwy7>h3*-DJEA?fvSppv$6gnTT&NY2}d}dsc<$-3_@3FBM031A*-EM8k2EZBbn~o*3 z?41F+1Pp+$2J24=_Xc&SySq2vWiaS8d_8s!!Ayf`{n;CuHVCi)eg(F`@S}cvIS-)x zusDqk`hmyokp1fR#!}eEPZd6W^8{GNb^vy3wxfGzKxYrWN-bgVEesO#)zktrvI;Jpc&t z-OGR4PUnhGh$637RVV@^P|?;LDkpvg45oljP6`qstM$3tO_I^lN^vh=y^PgoJ?Ywh zZlugxZ3KdV(tB3%Qd#<8LsMUo8?t8vZY;dHtS|D1ZnJ+ z&QlGnJPnPnjCR@wmz@FibOWs9IokS46YcNf>IY1SL!ThmjrG$f9k_3 z*18Qs{#geo_vy0$)y=gA;82r%m3_QedNMuLUzFclInVb76Z<}%H9ckdJ~vyIOpitn z15t(ct4fx0$>hl&yIKQ@2aRR#ob*hO!w=%6u?r#AKzlc>}LV#gDG&p&Odje1h=C45Aq5r~Cit9G^)d<<&ma+ #include #include @@ -13,15 +14,14 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 16; // feel free to change the size of array +const int SIZE = 1 << 26; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; int *c = new int[SIZE]; - -int main(int argc, char* argv[]) { +void scanTest() { // Scan tests - + std::vector timeVector; printf("\n"); printf("****************\n"); printf("** SCAN TESTS **\n"); @@ -37,21 +37,21 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, b); printDesc("cpu scan, power-of-two"); StreamCompaction::CPU::scan(SIZE, b, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(SIZE, b, true); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)"); + //printArray(SIZE, b, true); zeroArray(SIZE, c); printDesc("cpu scan, non-power-of-two"); StreamCompaction::CPU::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(NPOT, b, true); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)"); + //printArray(NPOT, b, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); - printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - printArray(SIZE, c, true); + printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); + //printArray(SIZE, c, true); printCmpResult(SIZE, b, c); /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan @@ -63,35 +63,35 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, c); printDesc("naive scan, non-power-of-two"); StreamCompaction::Naive::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - printArray(SIZE, c, true); + printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); + //printArray(SIZE, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, power-of-two"); StreamCompaction::Efficient::scan(SIZE, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); //printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, non-power-of-two"); StreamCompaction::Efficient::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); //printArray(NPOT, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); printDesc("thrust scan, power-of-two"); StreamCompaction::Thrust::scan(SIZE, c, a); - printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); //printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("thrust scan, non-power-of-two"); StreamCompaction::Thrust::scan(NPOT, c, a); - printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + printElapsedTime(StreamCompaction::Thrust::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); //printArray(NPOT, c, true); printCmpResult(NPOT, b, c); @@ -113,42 +113,51 @@ int main(int argc, char* argv[]) { zeroArray(SIZE, b); printDesc("cpu compact without scan, power-of-two"); count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)"); expectedCount = count; - printArray(count, b, true); + //printArray(count, b, true); printCmpLenResult(count, expectedCount, b, b); zeroArray(SIZE, c); printDesc("cpu compact without scan, non-power-of-two"); count = StreamCompaction::CPU::compactWithoutScan(NPOT, c, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)"); expectedNPOT = count; - printArray(count, c, true); + //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); zeroArray(SIZE, c); printDesc("cpu compact with scan"); count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); - printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)"); - printArray(count, c, true); + printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), timeVector, "(std::chrono Measured)"); + //printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); printDesc("work-efficient compact, power-of-two"); count = StreamCompaction::Efficient::compact(SIZE, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); //printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); printDesc("work-efficient compact, non-power-of-two"); count = StreamCompaction::Efficient::compact(NPOT, c, a); - printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); + printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), timeVector, "(CUDA Measured)"); //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); + + for (int i = 0; i < timeVector.size(); i++) + { + std::cout << timeVector[i] << std::endl; + } system("pause"); // stop Win32 console from closing on exit delete[] a; delete[] b; delete[] c; } + +int main(int argc, char* argv[]) { + scanTest(); +} diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index 025e94a..e5d63c2 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -5,7 +5,7 @@ #include #include #include - +#include template int cmpArrays(int n, T *a, T *b) { for (int i = 0; i < n; i++) { @@ -68,9 +68,10 @@ void printArray(int n, int *a, bool abridged = false) { } printf("]\n"); } - -template -void printElapsedTime(T time, std::string note = "") +// +//template +void printElapsedTime(float time, std::vector& timeArray, std::string note = "") { std::cout << " elapsed time: " << time << "ms " << note << std::endl; + timeArray.push_back(time); } diff --git a/stream_compaction/common.h b/stream_compaction/common.h index b76a9c9..119b5ac 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -13,7 +13,7 @@ #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) -#define BLOCK_SIZE 256 +#define BLOCK_SIZE 128 /** * Check for CUDA errors; print and exit if there was a problem. diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 818d30b..56daada 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -20,13 +20,14 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO - odata[0] = idata[0]; - for (size_t i = 1; i < n; i++) { - odata[i] = odata[i - 1] + idata[i]; + odata[0] = 0; + odata[1] = idata[0]; + for (size_t i = 2; i < n; i++) { + odata[i] = odata[i - 1] + idata[i - 1]; } timer().endCpuTimer(); } - void scanNoTimer(int n, int* odata, const int* idata) { + void scanInclusive(int n, int* odata, const int* idata) { // TODO odata[0] = idata[0]; for (size_t i = 1; i < n; i++) { @@ -48,7 +49,6 @@ namespace StreamCompaction { j++; } } - timer().endCpuTimer(); return j; } @@ -67,7 +67,7 @@ namespace StreamCompaction { { boolFlag[i] = idata[i] == 0 ? 0 : 1; } - scanNoTimer(n, scanRes, boolFlag); // odata: scan result + scanInclusive(n, scanRes, boolFlag); // odata: scan result for (size_t i = 0; i < n; i++) { diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..71a5528 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -2,6 +2,8 @@ #include #include "common.h" #include "efficient.h" +#include "common.cu" +using namespace StreamCompaction::Common; namespace StreamCompaction { namespace Efficient { @@ -11,16 +13,140 @@ namespace StreamCompaction { static PerformanceTimer timer; return timer; } + //Down sweep + __global__ void kernUpSweep(int* g_idata, int n, int d) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + int pow2dplus1 = 1 << d + 1; + int pow2d = 1 << d; // d0 p1, d1 p2 + if (index % pow2dplus1 == 0) { + g_idata[index + pow2dplus1 - 1] += g_idata[index + pow2d - 1]; // + } + //if (index % (1 << (d + 1)) == 0) { + // g_idata[index + (1 << (d + 1)) - 1] += g_idata[index + (1 << d) - 1]; + //} + } + __global__ void kernDownSweep(int* g_idata, int n, int d) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + + //if (index % (1 << (d + 1)) == 0) { + // int t = g_idata[index + (1 << d) - 1]; + // g_idata[index + (1 << d) - 1] = g_idata[index + (1 << (d + 1)) - 1]; + // g_idata[index + (1 << (d + 1)) - 1] += t; + //} + int pow2dplus1 = 1 << d + 1; + int pow2d = 1 << d; // d0 p1, d1 p2 + + if (index % pow2dplus1 == 0) { + int temp = g_idata[index + pow2d - 1]; + g_idata[index + pow2d - 1] = g_idata[index + pow2dplus1 - 1]; + g_idata[index + pow2dplus1 - 1] += temp; + } + } + __global__ void kernScan(const int* idata, int* odata, int n, int d) { + + } + + __global__ void kernSetZero(int n, int* idata) { + idata[n - 1] = 0; + } + + /*__global__ void kernUpSweep(int d, int n, int* idata) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + if (index % (1 << (d + 1)) == 0) { + idata[index + (1 << (d + 1)) - 1] += idata[index + (1 << d) - 1]; + } + } + + + __global__ void kernDownSweep(int d, int n, int* idata) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + if (index % (1 << (d + 1)) == 0) { + int t = idata[index + (1 << d) - 1]; + idata[index + (1 << d) - 1] = idata[index + (1 << (d + 1)) - 1]; + idata[index + (1 << (d + 1)) - 1] += t; + } + }*/ + + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + + // allocate memory + int paddedSize = 1 << ilog2ceil(n); + dim3 blocksPerGrid((paddedSize + BLOCK_SIZE - 1) / BLOCK_SIZE); + + int* dev_in; + cudaMalloc((void**)&dev_in, paddedSize * sizeof(int)); + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + timer().startGpuTimer(); - // TODO + + // Up Sweep Phase + for (int d = 0; d <= ilog2ceil(paddedSize) - 1; d++) { + kernUpSweep <<< blocksPerGrid, BLOCK_SIZE >>> ( dev_in, paddedSize, d); + checkCUDAError("kernUpSweep failed"); + + } + kernSetZero << < 1, 1 >> > (paddedSize, dev_in); + + // Down Sweep Phase + for (int d = ilog2ceil(paddedSize) - 1; d >= 0; d--) { + kernDownSweep <<< blocksPerGrid, BLOCK_SIZE >>> (dev_in, paddedSize, d); + checkCUDAError("kernDownSweep failed"); + + } timer().endGpuTimer(); + + // send the data to host + cudaMemcpy(odata, dev_in, sizeof(int) * (n), cudaMemcpyDeviceToHost); + cudaFree(dev_in); + } + void scanNoTimer(int n, int* odata, const int* idata) { + + // allocate memory + int paddedSize = 1 << ilog2ceil(n); + dim3 blocksPerGrid((paddedSize + BLOCK_SIZE - 1) / BLOCK_SIZE); + + int* dev_in; + cudaMalloc((void**)&dev_in, paddedSize * sizeof(int)); + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + + // Up Sweep Phase + for (int d = 0; d <= ilog2ceil(paddedSize) - 1; d++) { + kernUpSweep << < blocksPerGrid, BLOCK_SIZE >> > (dev_in, paddedSize, d); + checkCUDAError("kernUpSweep failed"); + + } + kernSetZero << < 1, 1 >> > (paddedSize, dev_in); + + // Down Sweep Phase + for (int d = ilog2ceil(paddedSize) - 1; d >= 0; d--) { + kernDownSweep << < blocksPerGrid, BLOCK_SIZE >> > (dev_in, paddedSize, d); + checkCUDAError("kernDownSweep failed"); + + } + + // send the data to host + cudaMemcpy(odata, dev_in, sizeof(int) * (n), cudaMemcpyDeviceToHost); + cudaFree(dev_in); + } /** * Performs stream compaction on idata, storing the result into odata. * All zeroes are discarded. @@ -31,10 +157,43 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { + dim3 blocksPerGrid((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + + int* count = new int[2]; + + int* dev_in; + int* dev_bool; + int* dev_ScanRes; + int* dev_out; + cudaMalloc((void**)&dev_in, n * sizeof(int)); + cudaMalloc((void**)&dev_bool, n * sizeof(int)); + cudaMalloc((void**)&dev_ScanRes, n * sizeof(int)); + cudaMalloc((void**)&dev_out, n * sizeof(int)); + + cudaMemcpy(dev_in, idata, sizeof(int) * n, cudaMemcpyHostToDevice); + timer().startGpuTimer(); // TODO + kernMapToBoolean << > > (n, dev_bool, dev_in); + scanNoTimer(n, dev_ScanRes, dev_bool); + kernScatter << > > (n, dev_out, dev_in, dev_bool, dev_ScanRes); timer().endGpuTimer(); - return -1; + + cudaMemcpy(count, &dev_bool[n - 1], sizeof(int), cudaMemcpyDeviceToHost); + //cudaMemcpy(count + 1, dev_ScanRes + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + + //size equals to last of boolean array and last of boolean prefix sum array + int size; + cudaMemcpy(&size, &dev_ScanRes[n - 1], sizeof(int), cudaMemcpyDeviceToHost); // copy the last element of scan result + size += count[0]; + + cudaMemcpy(odata, dev_out, sizeof(int) * size, cudaMemcpyDeviceToHost); + + cudaFree(dev_in); + cudaFree(dev_bool); + cudaFree(dev_ScanRes); + cudaFree(dev_out); + return size; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 09a4235..33bf0ec 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -18,7 +18,7 @@ namespace StreamCompaction { return; } - int pow2dminus1 = std::pow(2, d - 1); + int pow2dminus1 = 1 << d - 1; if (index >= pow2dminus1) { odata[index] = idata[index] + idata[index - pow2dminus1]; } @@ -26,6 +26,7 @@ namespace StreamCompaction { odata[index] = idata[index]; } } + // Part 5 //__global__ void scan(float* g_odata, float* g_idata, int n) { // extern __shared__ float temp[]; // allocated on invocation // int thid = threadIdx.x; @@ -45,6 +46,7 @@ namespace StreamCompaction { // g_odata[thid] = temp[pout * n + thid]; // write output //} + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ @@ -68,7 +70,9 @@ namespace StreamCompaction { timer().endGpuTimer(); // send the data to host - cudaMemcpy(odata, dev_in, sizeof(int) * (n), cudaMemcpyDeviceToHost); + // shift the output by 1 for exclusive scanc + odata[0] = 0; + cudaMemcpy(odata + 1, dev_in, sizeof(int) * (n - 1), cudaMemcpyDeviceToHost); cudaFree(dev_in); cudaFree(dev_out); } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..cec21c5 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -5,6 +5,7 @@ #include #include "common.h" #include "thrust.h" +#include namespace StreamCompaction { namespace Thrust { @@ -18,11 +19,20 @@ 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` // example: for device_vectors dv_in and dv_out: // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + + thrust::host_vector host_in(idata, idata + n); + thrust::device_vector dev_in = host_in; + thrust::device_vector dev_out(n); + + + timer().startGpuTimer(); + thrust::exclusive_scan(dev_in.begin(), dev_in.end(), dev_out.begin()); timer().endGpuTimer(); + thrust::copy(dev_out.begin(), dev_out.end(), odata); + } } }