From a90bd3fbf5e66320fe36985be92f33b9cce9d012 Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 15:54:58 -0400 Subject: [PATCH 01/12] efficient first commit --- stream_compaction/common.cu | 8 +++ stream_compaction/cpu.cu | 35 ++++++++++- stream_compaction/efficient.cu | 107 ++++++++++++++++++++++++++++++++- 3 files changed, 147 insertions(+), 3 deletions(-) diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..f314c9f 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -24,6 +24,9 @@ namespace StreamCompaction { */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { // TODO + int thid = threadIdx.x + (blockIdx.x*blockDim.x); + bools[thid] = (idata[thid] != 0); + } /** @@ -33,6 +36,11 @@ namespace StreamCompaction { __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { // TODO + int thid = threadIdx.x + blockIdx.x*blockDim.x; + if(bools[thid] == 1) + { + odata[indices[thid]] = idata[thid]; + } } } diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..ccca286 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -20,6 +20,12 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + for(int i = 0; i= 1; offset /= 2) // traverse down tree & build scan + { + + int ai = offset * (2 * thid + 1) - 1; + int bi = offset * (2 * thid + 2) - 1; + if ((ai < n) && (bi < n)) { + float t = idata[ai]; + idata[ai] = idata[bi]; + idata[bi] += t; + } + + __syncthreads(); + + } + if (thid < n/2) + odata[2 * thid] = idata[2 * thid]; // write results to device memory + __syncthreads(); + if(thid < n/2) + odata[2 * thid + 1] = idata[2 * thid + 1]; + __syncthreads(); + } + /** * 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 *g_odata; + int *g_idata; + cudaMalloc((void**)&g_odata,n*sizeof(int)); + + cudaMalloc((void**)&g_idata,n*sizeof(int)); + + + + cudaMemcpy(g_idata,idata,n*sizeof(int),cudaMemcpyHostToDevice); + + int threadsPerBlock = 256; + int blocksPerGrid = ((n/2) + threadsPerBlock - 1) / threadsPerBlock; + scan_global << > > (n, g_odata, g_idata); + + cudaMemcpy(odata,g_odata,sizeof(int)*n,cudaMemcpyDeviceToHost); + + cudaFree(g_idata); + cudaFree(g_odata); timer().endGpuTimer(); } @@ -33,8 +95,51 @@ namespace StreamCompaction { int compact(int n, int *odata, const int *idata) { timer().startGpuTimer(); // TODO + int *g_bools = 0; + int *bools = 0; + int *temp_array = (int*)malloc(sizeof(int)*n); + int* indices; + int* g_idata; + int *g_odata; + + int threadsPerBlock = 256; + int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock; + cudaError_t result = cudaMalloc((void**)(&g_bools), n * sizeof(int)); + if (result != cudaSuccess) { + fprintf(stderr, "Kernel launch 1 failed: %s\n", cudaGetErrorString(result)); + cudaFree(g_bools); + timer().endGpuTimer(); + return -1; + } + result = cudaMalloc((void**)(&g_idata), n * sizeof(int)); + if (result != cudaSuccess) { + fprintf(stderr, "Kernel launch 2 failed: %s\n", cudaGetErrorString(result)); + cudaFree(g_bools); + cudaFree(g_idata); + timer().endGpuTimer(); + return -1; + } + cudaMemcpy(g_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + cudaMalloc((void**)&g_odata,sizeof(int)*n); + StreamCompaction::Common::kernMapToBoolean<<>>(n, g_bools, g_idata); + bools = (int*)malloc(n*sizeof(int)); + cudaDeviceSynchronize(); + cudaMemcpy(bools,g_bools,n*sizeof(int),cudaMemcpyDeviceToHost); + + result = cudaMalloc(&indices, n * sizeof(int)); timer().endGpuTimer(); - return -1; + scan(n, temp_array, bools); + timer().startGpuTimer(); + cudaMemcpy(indices, temp_array, n * sizeof(int), cudaMemcpyHostToDevice); + + StreamCompaction::Common::kernScatter<<>>(n, g_odata,g_idata, g_bools, indices); + cudaMemcpy(odata,g_odata,n*sizeof(int),cudaMemcpyDeviceToHost); + cudaFree(bools); + cudaFree(indices); + cudaFree(g_odata); + timer().endGpuTimer(); + + return temp_array[n-1]; } } } From f3d99f5eac594772258375bfb2276d8a33395b97 Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 17:25:26 -0400 Subject: [PATCH 02/12] scan pass --- stream_compaction/naive.cu | 37 ++++++++++++++++++++++++++++++++++++- 1 file changed, 36 insertions(+), 1 deletion(-) diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..b3aa6af 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -3,6 +3,8 @@ #include "common.h" #include "naive.h" +#define block_size 256 + namespace StreamCompaction { namespace Naive { using StreamCompaction::Common::PerformanceTimer; @@ -12,13 +14,46 @@ namespace StreamCompaction { return timer; } // TODO: __global__ - + __global__ void scan_global(int n, int *odata, int *idata, int *temp) + { + int thid = threadIdx.x + (blockIdx.x * blockDim.x); + int pout = 0, pin = 1; + // Load input into global memory. + // This is exclusive scan, so shift right by one + // and set first element to 0 + + temp[pout * n + thid] = (thid > 0) ? 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] + temp[pin* n + thid]; + else + temp[pout * n + thid] = temp[pin * n + thid]; + __syncthreads(); + } + 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) { timer().startGpuTimer(); // TODO + int *g_odata,*g_idata,*temp; + cudaError_t result = cudaMalloc((void**)&g_idata, n * sizeof(int)); + result = cudaMalloc((void**)&g_odata,n*sizeof(int)); + result = cudaMalloc((void**)&temp,2 * n * sizeof(int)); + cudaMemcpy(g_idata,idata,sizeof(int)*n,cudaMemcpyHostToDevice); + int threadsPerBlock = 256; + int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock; + scan_global<<>>(n,g_odata,g_idata,temp); + cudaMemcpy(odata,g_odata,sizeof(int)*n,cudaMemcpyDeviceToHost); + for (int i = 0; i < n; i++) + { + printf("%d %d\n", idata[i], odata[i]); + } timer().endGpuTimer(); } } From 99ca572ecb62fb35401ec93376fb01681886288a Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 17:37:11 -0400 Subject: [PATCH 03/12] zero padding for non-two-power array size --- stream_compaction/naive.cu | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index b3aa6af..139f7d8 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -42,17 +42,18 @@ namespace StreamCompaction { timer().startGpuTimer(); // TODO int *g_odata,*g_idata,*temp; - cudaError_t result = cudaMalloc((void**)&g_idata, n * sizeof(int)); - result = cudaMalloc((void**)&g_odata,n*sizeof(int)); - result = cudaMalloc((void**)&temp,2 * n * sizeof(int)); + int zeropadded_n = pow(2, ilog2ceil(n)); + cudaError_t result = cudaMalloc((void**)&g_idata, zeropadded_n * sizeof(int)); + result = cudaMalloc((void**)&g_odata,zeropadded_n*sizeof(int)); + result = cudaMalloc((void**)&temp,2 * zeropadded_n * sizeof(int)); cudaMemcpy(g_idata,idata,sizeof(int)*n,cudaMemcpyHostToDevice); int threadsPerBlock = 256; - int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock; - scan_global<<>>(n,g_odata,g_idata,temp); + int blocksPerGrid = (zeropadded_n + threadsPerBlock - 1) / threadsPerBlock; + scan_global<<>>(zeropadded_n,g_odata,g_idata,temp); cudaMemcpy(odata,g_odata,sizeof(int)*n,cudaMemcpyDeviceToHost); for (int i = 0; i < n; i++) { - printf("%d %d\n", idata[i], odata[i]); + //printf("%d %d\n", idata[i], odata[i]); } timer().endGpuTimer(); } From c0b031cefd8aeedd8b5b588562707a075a1aed6f Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 18:13:14 -0400 Subject: [PATCH 04/12] non-power for efficient compact --- stream_compaction/efficient.cu | 50 ++++++++++++++++++---------------- 1 file changed, 26 insertions(+), 24 deletions(-) diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 97768dc..e019f99 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -64,17 +64,16 @@ namespace StreamCompaction { // TODO int *g_odata; int *g_idata; - cudaMalloc((void**)&g_odata,n*sizeof(int)); + int zeropadded_n = pow(2, ilog2ceil(n)); + cudaMalloc((void**)&g_odata, zeropadded_n * sizeof(int)); - cudaMalloc((void**)&g_idata,n*sizeof(int)); + cudaMalloc((void**)&g_idata,zeropadded_n * sizeof(int)); - - cudaMemcpy(g_idata,idata,n*sizeof(int),cudaMemcpyHostToDevice); int threadsPerBlock = 256; - int blocksPerGrid = ((n/2) + threadsPerBlock - 1) / threadsPerBlock; - scan_global << > > (n, g_odata, g_idata); + int blocksPerGrid = ((zeropadded_n /2) + threadsPerBlock - 1) / threadsPerBlock; + scan_global << > > (zeropadded_n, g_odata, g_idata); cudaMemcpy(odata,g_odata,sizeof(int)*n,cudaMemcpyDeviceToHost); @@ -97,49 +96,52 @@ namespace StreamCompaction { // TODO int *g_bools = 0; int *bools = 0; - int *temp_array = (int*)malloc(sizeof(int)*n); int* indices; int* g_idata; int *g_odata; - + int zeropadded_n = pow(2,ilog2ceil(n)); + printf("zeropadded = %d\n", zeropadded_n); int threadsPerBlock = 256; - int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock; - cudaError_t result = cudaMalloc((void**)(&g_bools), n * sizeof(int)); + int blocksPerGrid = (zeropadded_n + threadsPerBlock - 1) / threadsPerBlock; + int *temp_array = (int*)malloc(sizeof(int)*zeropadded_n); + cudaError_t result = cudaMalloc((void**)(&g_bools), zeropadded_n * sizeof(int)); if (result != cudaSuccess) { - fprintf(stderr, "Kernel launch 1 failed: %s\n", cudaGetErrorString(result)); + fprintf(stderr, "Mem alloc failed: %s\n", cudaGetErrorString(result)); cudaFree(g_bools); timer().endGpuTimer(); return -1; } - result = cudaMalloc((void**)(&g_idata), n * sizeof(int)); + result = cudaMalloc((void**)(&g_idata), zeropadded_n * sizeof(int)); + cudaMemset(g_idata, 0, zeropadded_n*sizeof(int)); if (result != cudaSuccess) { - fprintf(stderr, "Kernel launch 2 failed: %s\n", cudaGetErrorString(result)); + fprintf(stderr, "Mem alloc failed: %s\n", cudaGetErrorString(result)); cudaFree(g_bools); cudaFree(g_idata); timer().endGpuTimer(); return -1; } cudaMemcpy(g_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); - cudaMalloc((void**)&g_odata,sizeof(int)*n); - StreamCompaction::Common::kernMapToBoolean<<>>(n, g_bools, g_idata); - bools = (int*)malloc(n*sizeof(int)); + cudaMalloc((void**)&g_odata,sizeof(int)* zeropadded_n); + StreamCompaction::Common::kernMapToBoolean<<>>(zeropadded_n, g_bools, g_idata); + bools = (int*)malloc(zeropadded_n*sizeof(int)); cudaDeviceSynchronize(); - cudaMemcpy(bools,g_bools,n*sizeof(int),cudaMemcpyDeviceToHost); - - result = cudaMalloc(&indices, n * sizeof(int)); + cudaMemcpy(bools,g_bools,zeropadded_n*sizeof(int),cudaMemcpyDeviceToHost); + + result = cudaMalloc(&indices, zeropadded_n * sizeof(int)); timer().endGpuTimer(); - scan(n, temp_array, bools); + scan(zeropadded_n, temp_array, bools); + timer().startGpuTimer(); - cudaMemcpy(indices, temp_array, n * sizeof(int), cudaMemcpyHostToDevice); + cudaMemcpy(indices, temp_array, zeropadded_n * sizeof(int), cudaMemcpyHostToDevice); - StreamCompaction::Common::kernScatter<<>>(n, g_odata,g_idata, g_bools, indices); - cudaMemcpy(odata,g_odata,n*sizeof(int),cudaMemcpyDeviceToHost); + StreamCompaction::Common::kernScatter<<>>(zeropadded_n, g_odata,g_idata, g_bools, indices); + cudaMemcpy(odata,g_odata,zeropadded_n*sizeof(int),cudaMemcpyDeviceToHost); cudaFree(bools); cudaFree(indices); cudaFree(g_odata); timer().endGpuTimer(); - return temp_array[n-1]; + return temp_array[zeropadded_n-1]; } } } From c7f8dea54adc99796b5c62a0f95e66fb5fa7f5f0 Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 18:48:18 -0400 Subject: [PATCH 05/12] thrust implementation --- stream_compaction/thrust.cu | 10 +++++++++- 1 file changed, 9 insertions(+), 1 deletion(-) diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..d907c7f 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -21,7 +21,15 @@ namespace StreamCompaction { 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 thrust_idata(idata, idata + n); + thrust::host_vector thrust_odata(odata, odata + n); + + thrust::device_vector thrust_dev_idata = thrust_idata; + thrust::device_vector thrust_dev_odata = thrust_odata; + + thrust::exclusive_scan(thrust_dev_idata.begin(), thrust_dev_idata.end(), thrust_dev_odata.begin()); + thrust::copy(thrust_dev_odata.begin(), thrust_dev_odata.end(), thrust_odata.begin()); + std::memcpy(odata, thrust_odata.data(), n * sizeof(int)); timer().endGpuTimer(); } } From cef2591c4dc3d0c8bc7c55ca73928959a967976c Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 18:57:28 -0400 Subject: [PATCH 06/12] info update --- README.md | 12 +++--------- 1 file changed, 3 insertions(+), 9 deletions(-) diff --git a/README.md b/README.md index 0e38ddb..56859b6 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,6 @@ 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) - -### (TODO: Your README) - -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) - +* MANVI AGARWAL + * [LinkedIn](https://www.linkedin.com/in/manviagarwal27/) +* Tested on: Windows 11, AMD Ryzen 5 7640HS @ 4.30GHz 16GB, GeForce RTX 4060 8GB(personal) From 860ee9d6292b4d99a02c5ce3aaa88aef96a41a68 Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 21:53:06 -0400 Subject: [PATCH 07/12] Readme Update --- README.md | 82 ++++++++++++++++++++++++++++++++++++++++ img/ComparsionChart.png | Bin 0 -> 22277 bytes 2 files changed, 82 insertions(+) create mode 100644 img/ComparsionChart.png diff --git a/README.md b/README.md index 56859b6..5b207b0 100644 --- a/README.md +++ b/README.md @@ -6,3 +6,85 @@ CUDA Stream Compaction * MANVI AGARWAL * [LinkedIn](https://www.linkedin.com/in/manviagarwal27/) * Tested on: Windows 11, AMD Ryzen 5 7640HS @ 4.30GHz 16GB, GeForce RTX 4060 8GB(personal) + +### Performance Analysis + +This repository compares the implementation of scan or prefix scan algorithm on CPU and GPU. The code in `main.cpp` does correctness check as well as profiles individual implementations to compute the time. The output for the test code is: + +``` +**************** +** SCAN TESTS ** +**************** + [ 31 35 17 38 8 45 8 15 34 30 36 13 29 ... 7 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 0.0012ms (std::chrono Measured) + [ 0 31 66 83 121 129 174 182 197 231 261 297 310 ... 6271 6278 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 0.0013ms (std::chrono Measured) + [ 0 31 66 83 121 129 174 182 197 231 261 297 310 ... 6206 6242 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 0.640896ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== + elapsed time: 0.16288ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 0.245184ms (CUDA Measured) + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 0.115168ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 16.073ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 1.03667ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 3 1 1 2 0 3 2 1 0 0 0 1 3 ... 1 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 0.0012ms (std::chrono Measured) + [ 3 1 1 2 3 2 1 1 3 3 2 3 3 ... 3 1 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 0.0011ms (std::chrono Measured) + [ 3 1 1 2 3 2 1 1 3 3 2 3 3 ... 2 2 ] + passed +==== cpu compact with scan ==== + elapsed time: 0.0055ms (std::chrono Measured) + [ 3 1 1 2 3 2 1 1 3 3 2 3 3 ... 3 1 ] + passed +==== work-efficient compact, power-of-two ==== +zeropadded = 256 + elapsed time: 0.077728ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== +zeropadded = 256 + elapsed time: 0.083264ms (CUDA Measured) + passed +``` + +## Computation comparison + +![](img/ComparisonChart.png) + +## Insights + +The results favor CPU implementation heavily in terms of timing. There are a couple of reasons for GPU implementation to be much slower than CPU one. All these limitations can be explored to look for optimization for GPU implementation. + +Following are some of the reasons for GPU implementation to be much slower: + +**1. Expensive memory operations:** + For the implementation, I've used global memory and before calling GPU, the data is transferred from CPU memory to global memory which adds to the latency of GPU implementation. + + **2. Small computations and small data set:** + Since each thread doesn't have a lot of computation, the time it takes to transfer data from host to device and back ends up surpassing the reduction in time that parallelization of computation attempts to achieve. + +**3. Warp Partitioning:** +With each iteration, number of threads working reduce but the corresponding warps remain active. This causes divergent wraps and hence the GPU is under-utilized. + + diff --git a/img/ComparsionChart.png b/img/ComparsionChart.png new file mode 100644 index 0000000000000000000000000000000000000000..a9bf25551382d89434072d8ed79cfe62b1bfaae8 GIT binary patch literal 22277 zcmeIa2UJtryDu8I1(hwRC>Fq2QIsN54Pfl(R*F=S5~L$lLO=w9%0@s0HX@yXf(Rt^ zPz;eEDxCxs1B4cp5(!BVB7u;E+!Z2w|L+_BckUVQj&a9*XB@{+aILxKH|JM=Uz=;i zni?B!+PHfo2n5=6`OigMGsCl>GLrlx@SipAXN=B( zK$VG-{F|biWu{jxF7)^J_x1HrD3pPL0pQ<%{j=Tifr?+es(*qyJn2N>6VsSH8yg!3 z2M1?oXScLcPft%@Utc&Je&@~|BoY}J85t836CWS{l`C}!djwrYj@*kIQ!qs3!PBz zZh`05f7MFVe>_b^9P&D?U9NmuXX{!1xSEhDTnlpiUC4~!bqXh=K_CxaSA%v+lUIQrN7F%|mUDn8MxR4K_hw8LKyU68Nr4>o|H}-6AGPA9 zfxq^BF*bZqWGqjpYLMkfq`Cz<8y0@K=oZFOE4ouAYgBJB9T*C>d>MyHi_cx-kq5c0 zYk|_y3#)~O{#InXzCQRa2I6!V&yiVWU~yO~8V&c0uH(3yfXExp%o|S^8K)9Q1Qw?^ zkz?OIr}w9La~vZ$YfMqhTQ}`^lIIvp4_K%X8l zhknb4e4H#xkhiieb-_e>!&JM8CThNz)egbs#rkR!oZfCQN1vNBo;sMjIV9!SG=DAp zVb?l6BXIH$Xx4rNj$hynV=Iv9)VRj>nJza<90#tzF==f=xe_ONFd=vg~Yano_$a2 z=5PD>k{revxoc_60i#?8Xt09hN8VmWsQ(=F$2#S74){SXz>gk#)I>_a8LOS$tXRdc zqdB+ws&i*+a!HDvq{#dc=KJ1jS2@>NPDCy}hf50UzFXRG@tCSqE)4f}>7ic5A9iGG zGZcA}yLue`?8?=$Cvst>NSABMIw!y3${-kGq5X*{wspi@j%cwz9r!cE-U}%z`>dH+ z`BEryWiYEijyLHQe54Ui`ku9R#xD2N;N-dK$cpJEvha;uh2ucPdgMlOEa?Nit9H;f zw6yMc_ZMPKqociM-^uvxIX2^nvLpuNdC1n(_;4SVHCOsX87|nmF3^(pid$oy9uMNd z{2T_uKUty(P}8#Y_3y_v!Xt^h_<=nfE1K6txzYm4<{YMHq3>}^y>(f?j-dhsT0vM&^h@gu#h)v^9<~}|UQ$JibCRrBl6~CZ)Gd$jpc_@01`KRYhD3S3~^%FtDTDW@p#{&xY+z)t>ao{>@L#uYID`8;5cU<4L zo&C{)jr6NCajQ|-$T*XR4(0Av&@oJh<-t_PwClz25EGI;4-}cth>I4`MYah z^C#F~0TffaeJl^q5*0vq8^29o+ht5_v~l^Qt10QbXnAt@VwazyoFJ|m&|cmfxhlsS zfc9<$wD$-7YNz|hmIeZ*;q73*VDnP=&3C7KFqUZB)i0VRHt(L|4W9opFj`~Em)*wQAh zTCk(zsC-F+K!Oy9@TxQ}_cc+SBVM~=jR@e$o%%D#d%u++PQty30E6klvn9Q__aY|V2L6-JKOMRFs0$;(9Wrt7<;{5 z?E}n~OSlCo%T{AQ#&Niu#&}XV$E;hvCeXs&*}AWJUBP4G%V&5;h0L|V#o7|SxlSK5 zR>xVi0B{@L^H3#y!Zv8(vc+LYB*q+eXwd5FEsSYh?1`BjL-D;1u8TFkeOo+&j7;?| zQCI`BUh8%l-!fSWhhX4c{CcX#Yd`0R0z1z-4E($FQ^_-F>2^groocVDycD*0!gJey z?vS!-C7HsfKKr6l@ZZPIQbnt3ZGf7|UvW~?JcUPi$TGH~-CcQ#xhYB`f1Iqi-N1?RcdlVktjEapO(VvY zYaOkf$(p|T5vyL-}kM_lEl#VGJ1!=b~Sob)_q2L z4?1Bulx)W_39!V)8LR#zUh6bQ5Qbj08~I4O;%o4qosu&r`)n10Y~7oO_w87mC~6)tY@VMQEARB; z(fT;&DOZpT`Si8(O=SZ2+j{?xm7KLKB>aK(s_q?J0{Fphr?*XY-12)ktN~*bRED0f ziW3}@oz^3JmORiyzH68^E2+6bKDi*#p$l;1HhUctTxz9IXZQhjq@Nsio7VQqFX0B` zy**jokCO+sjdp8duC95f8Z9fr3ip9v+24`d$gv3*5sZ)K#BOB$j<@zE&iJ^}JCZI3 z$`5uTz77z|G);uGsy|Q`=uxSP*&Y8AH3Y_4xHy5|fh@2evOK}c%1>p{hb#tVI91EE zRIwQfK_{z4tTyN&V8yJ~L_7PzRi_zNjI!A>dVk^VnXd++&mrIU&f{-(vH6uV{Ua&)MTjf=Nkt#SK6)V78Xs8p z#KqE-;avx=EK*pPGC21qeCcyTzI`!-b=pxV6QpD*;_iSXlF{!)XnfCD3k`Z`dAPBr z{N1r;W%;6n;nvTmvi=ktj$5tlG6Z0q9ywz~6BgNjOwbR}lnlDdY^Y6rC;gTn(Kq5r zwK=|U4KT~y}iG^Z?!n-gAr?xl|J^ap?! zq-PR00j&!CTdM|YNjc0XQdQc4JQ|Y~H<7SoX6`ym5i$iSd)9Cf3!Q9DXmFe62MJn& zoXgqJk2R$|^Aa6D0dc?70n903IwpqN9LO~*IW8A5d!OEKem=0wE)MtfHh?+R)AEv*&Fkh2e*g>YyHoL!)2NJmou%k)i#6@M-W8~#INK^eE-<`XElsJsPi}3 z#K_R(JZEs$ABf#x3D456QcL$w7yx!VNt6RH1Fz`)`mP_p9T07!Y{IL7rqdw z_V@yTI&1YB!heU47fIagj>KrKY_;{&$>5xW#A8`$e*(OOd`4P02{9iOs^W~KF4lm% zhK8H0{9!fZ%AWHDkod&NYBcdr&+nCiW^?NqFe8GVGfCeIikS(qP8Aik{NnTPWXkSR=uzpPzes}x__{a+FUI;<3n)BxCgj9Y4l#+;Q(c)~bn7~hQPXM2?>^3BpXF2eYmm>c34%u@P--za@$uF*Ul?S}itB}Z2#B0HEV`l#1- zIBC|5be6Z{hVY>PT}40kdec0SIIey>;CXOD=D_?%DZm68&SAe0j7kY{MAmkdvptQ! z{Nyk22AlxL*Merwo{DUckf#JAspIED<`lexdg(e&u*Saf9|K;RxI07qfyuduFdyrV zaY@g*adTU;{KD$-vmXvA?#R^XztK8KEob^x0lWE4=x^wxC2VS0Qqmhv6#QINh43B` z20lH#(g6Sbc||nYYXO&3TRBSD{<e%osc!(RtX}sA;KY|aEyOZm?>T_yL1Eba0>mLC|V!w>kx0eka80%4OqI-RlQvQ|c zd|w+0z5pck7N{Cqzi!{ua_KAh!w0iI&p)jLa)FEN_1cZ*-<`PGjjF?t0>sb1WQrg` zg`Wc}uP?Cuh+d(ehYYnJJDL0;1be9QDxJ%@OIxvNf_~4^Q60|Qt8v*py?vxH)}ij< zs?b_PSjw`!2{%E{^)wS_JC<&x+V?ta>p#Zt-vC3c#(oJd&UmHD+JU?)C@mRT;bViC zk5AF&`dY`5%N({bRJHRL+?2bSXNxgG-k;Z?6ZG_pHU0RnZ9T~x>ip7u&c+oF3H&lJ zb!YHur{Vi@m4u}H)+wBr7oXAC;WCy$_OZ<*Q*cqco9Q&AdvuXl&Ww+X z?u8V5aLA|Ra-3TodVut&o;z(l6R*i0uA#LKzkMB8h6Wt!z!bl~vM($mmIC}jv2QKU zRb|HcgSTKAkNU=}gXF60Q5?0l3a;O0H;-P21VkLq21U|vSZRElA)0>Tv(+s zQ&2owI$IZQhxiH(^pLgUPC2*Ezu&%+tL|ue(6re?;u;#gdan}jGY@h)uShEIJmTK@ zYYOLM>7dv=NX%C9s;w6gn-8H^AACgL`-g?ZrETOhPm84fI)aWoa!q0HwSOKw2zyG6 zxN8PF3_SW@yoR|_tv!sUk39tc1~S`bg6Iv#cBK!O3bt#Vc;0RT0%bq-K(p274cbWr zzU6jkLCf;B>uBtg?^nyq3n!rNdbWDnnwyZnfxf(WZ-&w{e}Zpugt8CNkn2IZN=X*( zTJ2TGbnILxF5?e}$J}9pogja+{l%nw%G((pnb7s)<0gD%@oLcFV^8Smo3g78vIDS% zGFcFifBpXA1f|PdHZY0W!9bCnq)B8s~H-<0^#lD57$w10R^u>06Pp z?^*MllrLnvJM*LVZszW_GtQtlv9bA(Zbi*pq_@mZMysBSdExx^)u7YwzZDyYTzt?? zpjQC23zc)|PgRn1LO7PuMhYNZ&l!tWX&wg7cu4|J_^k$caNm{tNW}R|+nOK}Mm_^I z?Pp3+^q=xZ{y*^$proyd@<5C^PR_PX>cGDe6SG{*-6vZ?_hhyiBE~9ePH0wrHQ#e_ z*ijp?X!83;P>iiSnUenDshoQ7OBH1Q!7PXbVa+Pg<{Diz)}XU03hdkMHM&k40s8X! zf*C4QW?Mx1n|i{qdFx9lQ1Oo+%n_s+f{DV_0$d-tNLYHxW1Ol*yH?S3;Y0m~9fl z=xP1qw&(gpg1h$7!kwq8M|sSq;}N7P#wTg!Lz9^xY0#JR&E_a32k!Y93a%%03cLop z^{GNjKhV9N&E`PCzc~NDXu*HQSbnuO?rmP831d2iv+T9+AJRel_)uZ6-4YZEJMCR0 zwfQ&ls)x@Zr`KCZ>=fNJRdC#iUj653#J!vU*WJuRla6*XRw@ zT^i0y6*Jb=4Js~ogHm{%r_YP-Xp*_QSR-{26l|K3gr4k_^@>o~ya~PfcbO{)tE%T3 z4TPLbVK&0THC(mfWTSmoSA^=kkJKt?6dm+V1P(IE_?Bmr9d_XdvQX!yK4ecm7Y&O0FqT zFdTa8;9>allytl@{^vQ7PuQuLKnDuOJ+A=+S5%+FAQ{YQW>BiILxF=S*VRxe=@L7qVc3t>1Ey1dyw6#%r)? z|Ij<2S7F^ecrxIR>>bckW76!pRRpnqFFb_msX1<-PR%Oa7Xxj$d6yEoH%VHovK5!f61#D z1XDtE-cddp8iDFXY*r?>AGXnMK!)IQGD~+ytjStar2A~~LW@p0myOEa6)jV_Pg}Ii zRC@ntLBqGKps`53@!&nLDeMZHp0;^8!R=Smah?%OnEOsb7%!0*L~9!y8C&Q=AsdJb`KGq4;$G0#2&)WSky z{e3{+OcnO7`R|+mHJM-x5&3~G;jBjL`elg!v}G_nV*qtoqzT55q69!K0O0?01`!Cv znQPF=Qt13+XnM(P*m(3C&*zZIXOMb8u_kWj>|#g9H|E_`>e(n6aR>m_<&l~6ux#!N z7kAhax?T_Z5eQ|a$d@I^AGiN2-j)hdv~;v}_K?p6E|V;~jy^DWA~^m`&;HIkof2E} zm#F|T*R3b6*1jH#Cn%e`LQ=U~HXjfZc|zCERFPi+ShpM^e|^jrF`3JVu5~0FIP=Z6 zECVf@pb5j%kHsV|p*4ydGT_Os=1}F*lYfG4i-~+IN?=%q&W$pYaUYP*t^7AN;NR5q z#Jdy)Xvw@8&T@eEfO~(CN%z6MazHL%?xmZe0!TRnxtYTr50+`)TW_~i1vPBYR_n688Oyz5x-1Q_X~cm2=lsXzrI z)5Dg#7fDj&rb{Afq-)02kbsK7kg;kD8ZlANsNOiE#|OLa;!JSkgP7(}H=9o2xGA|!T$DC*36!%!)wBIpdWYy>EowX z?ze4Ov$TE{_q50DvYH*Rw8X^;;)PTqKzHqm++FGh-sjAkFOEXj)MvO;q6U9^(oQ%X z5b)r6bj2!XvEn)vCuHAtyvv&n(9;SB;>z}!&r2U_gHN()1R*`XB0HR7B&Z|QNy>9rI`lAE%mSA_%wVs&G z;uV=7Y6rUuJ}!iLj4W)doMh;;@2+;f`ex7SUdg zxkOfT;Xg3SI8mIotcA;{EmEofp6dSv!2bWot2ASTl}YDdvs0=oC9Lw5?Wu6?)P0L< z%kthx=DvdJ;5;*_gffB4-V4jxWBr{THqs_*To~sr)Hr=t3?UWdkjB}lcDJ}@asci% zOj#BlkwWRM=Ed%e#`OoImeqFSb(Cu=s&V0VUE-8-7Ek@JZDPfCL8!u^1JhFpJ_-k8S4Q46 z+~}{r>dE1bLQER)$p2JhPxbBNBdo z;y(ueHX|PS*4o|KZ1imKo%U2>$RI)G} z$-k{Sw4Jm0-bfvS&?WtPamI@C9MMq8R&Rd^vyiwgmeMvLaGjxK44B0D)}08Vxegkd znIQ!=*&m;c-m4?pq#EGCbA6y+{MG~p*U?~jWo^H%S5OECn|tb7*|DXkWW+p+xkpqG zTQ-ZMWOX^k5CtX`ARj@0>k&>!TJPk`Jw~&vk6zpat2m@fOg=O2N-{EA>DHAe6SKT*$Vs~F`VMvH9Efqf2eFJ3IA}0+1 z4W9zNZ*I5@YDyL~J$K<6WBFfjqYzp~!Zw#0Ou82f8M!OZY(s0bSFxwV9VK*aR*XMv z8o2K%p7X7>A!%tc`c1i*C}8+H-Afy?HJA1C%GYp|UyeXV*LC^Pw6h38ILG8m;EGeQ zfpbjwVoU+<)xy|}CyXz_;%G(_eJo@cgegtn5PwC5hG1XNVY&Nu1;9t&v)Xm%roYVm#9_u!` z3*{%P2ll-RN%+!}2Nb!N1z!O%Vaczl8-|r(k%X!qHNGoQDHlJGzk%SZsI8pbJnZa1 zrEs`m?zI8$Pg3U}YKqyg4v+hR?#QiVh05B{yR7|*Vn^KO$EtU%LNFF0ro&@B)Jf`A#jRzgcl zNSw%N@)@AgEXwPm=vT2gJ;!>W;ge(}SbynnE5j(`gYs`6d@=j}+P~twW|eL0ua!dP zZ8)sx(~p)XDEmR?wpGdLpe<>TPAla47ctAUHo{PdR21%-C2CGN>(!RO6n`nJ#{&Sy z2hP}_;whhTRR@5Ug-C*cX70f%hdw!T4He8iJ7stAGyIMv&%F8r*s(z#T?3IH6r>0Z?X=ry1Wu*unD2Qzo8@F zub3Y=_>XL@cahx z*7G5QVNZq^zFxBmAPzeHmD;&Wj5&nmjlJNwk%+#vy&(=hr!ss{ zlF&6yVhT#H%h~bR2qby_nE_eH8*x5u?NUX5of-AGDpQpM{J2%zrxAEhYAf6dRGwqib4n`O0Q{<@`D|* zM5DervNR!zNecgoNuZ+iU!Sa$xVhg}VP?*wLu1$LMCM9u_RM=xZPx#FKdZBCx(mu5 z2&w?;7liVjP|c>|yw1>#V&5Gg;0~ zMYg6cz*kpr_+RS>sVaQ`P=Z`ar*i0zbEEZ~ccUKSB^Ur3m<=(??%Vh9Bj zy&JDGl%NGlNW2@+m6F7-JxB#Z=GQJB?;`>0*)R5jYfDHSelOC6fx15~cyG;`WOjBk zKaaJ7n-p8kc}CgYvAY}$#<)F3h(nvD#is3uaM4K*n~8V1+4sC$_*4|sgl-)fdm!C4 z)n#0z*pMq&h5SvdixrS&o{H)wpJ*WPg$6_KIE2@cE-?2#m}~H6#n=Xw@B5WGw9O-Y z-F0SFA47MY>Ac5GHsRu0Xy@3A+yWi$Q~>u0Sa9aLpGnd`i6kYcili7>1d^HUc4__LPCG+>h1bZ2~S~PpYEdETcTZdqcy5`byznlsw<3 zhg}U$#IR<5^#lQr0`vq|Jc>}%^8=T!Tw9y z3>l9O-aVpdc+nki19TtpFn|*5A4>V%+hmqcGE_m9Pcqog2LuM*tMu1A1JTVqZ7n{{ zklTdotcnF|x1HPo1d-pwwW;MY^|G6B2*;cm{&b#}iAC-=6GP}hGEQeuhj3)Jbe#fi z&tHBOggNKEkd=E$2iEh`=xy3WQ9?YzcZ3fQ@87%i@4LgdWI{VkCQWA4KVX&d6}G3d zuE2e|Z36D)ii6x{Glf{0gM91XO;^F5$FE3L9ufcsHK)(Vfpo_wo`Q$Kj z6NQFw==C_ks=d3#6g0^cLU*_+db(FVOV`&uE6WV1YwQy1RSw;|j{E5dRQYH?#Pu^H zz4N7hZ=}Tmq!U`ozkXS;|E3i8-_W<5=>K;$UQ`|Xze*V{SfVU*%K4jhfwxFSb%cL8 z0TOse8~S6s_4HLSTmgr{maDU&_*rfB5cWF~8~=%T37)urJLfXgoSaIcuc6&O)N{@25mbN4(NwuvO~xtnlv}--<|S z#qAidUv5{F0a&&G5+g2rZ#P7|P{e`#V8pd}S8N=8ua$V=0LsIq!WJuyIdsY3Qn z5#Yd|%9JL^E8Fimsp6>wecclnh|uJ?=cKS{qnU| zR)*SP9q)>=63V2GZ%t6VBw;A?{ZT$d&OEsQ+?lrLD!(HBU$#lf*$xYI%yy^Xz_r)@ z-scO*yc@4&TEVWbF7!r`l+%Ufc;bSYLVK*xOo5w2m)5astR}TyL~K5`;%?h~Yc6wL z*ERP|bzvEKx5L(u&yzyUFT<23p=f_MZ3$k9lQXnE3#@A+XBx;))Hx>_H!bW`}g#qdN5I>5$=% zxewR=D&VPGO6b}zKm1R8^WV_(e;xGvcQh}8$$wPT7!3i3unbSThz+Nwm_#EU)KH}Y zHCx4HZSHHgptWlo4spGJqesh)nnE*_^(2uoadip%vv1X}CXew8NR~~ikqdR8Z^(uA z&VIvtamWox0uFoh>JNTwJ=j^sSpTc;f!I$z*$tdL9mxQd3U>ZF1c2C0=AI}B`7t{5 z7&x{2&&3en;0HFtf2Z6xaqf<~h)A{3o_ojS?C(_l`v?K>G_Ji02llPA2F{H{pH}|0 zp&}A6ivL0J2zw=<26j}a`=<}S0YG>$tRha3ah!8JMQQTN;f>^CQapH%H455AT?hIw z8x8c0f=-Z&rEG!YrkmpeG3E{B@UD+*L}zgi-9TX$lX_l4$kD+dP~jNRSggHfczgozf@;^1o$V02yQke9MD3aI8TfW6~s0UPU^nrX|V~)`1Q){scaqp&Rml z{_7aeCcO>_g`TaTBGTl-^$|7MxWPiT2peU|a-4mg4R|6z5KXv|HsfxPb?AnQU%vYW zsB2FmAZXi#&dArwV1C;TsA{=v9IpWP{7F~j+|zQt`9bv+bvw!{JCSW5f$0KPFhW)GMW;OL2B#l5PU$ZLRE8zz5rjj z==l+pv0awj(F!}D*w=hxoMMMAzdPIO%D0iko|o$uv-ZBm9&T2bjC?m8;>&)fzaaSrHKUfx`N*(9h=>?$=Ps{?6dgsM)BBM#fc#p%e?ltdg0- zbJlo+D+~KNyN;2ujJL2o-w1@?LbB3r)$F$qUZ10@*JaNjhaH`*c2UMo z2?r(;C$_*!ZMuFa>g1$y$WKBY%@lls;$abw=2gC$@bYRaq2u?6qBE>;zdr(H{iX-I zX@eQN7Gv7akYKg=YPloZ*LtKIz?ubr@fn z5&f++FBmsa+T7&`9O=RK*Z}-N?oaoCO-r5gvS@u>qC0;Mza?{chiwJNhEbR2g~h99 z1eBk@i+0|lT{v6+fc}2EM}J#4Cb%3*yy75yrX`IFICLNpicgHEduVv^t7fb|9P;kH zoILxIh@oDtK~<;5`I;D&LI$NZbbJvu5vSL~aYYFsVf@*66zlK^9_($`?JE&$^KPfk_}6uDQSw1|>aJU;y`aiU!czGiKfREa zw)W}UFNd`AuhbYFeK%A)hz8)!A-X?;#1DocbaS4K#(rTP{W`6+?}{*_@<=W`>u|`_ zIKN*11n$)BhNpt8IQ;BLkPNjP0KVH(Ec{gS*&yrm*8AaWO1uPVA?U`Eir|7!J1zIV zp?gEg;be8S%f@4OenZO`6klT3uq`Wp(=Z2C9UpPrT*1c-6(4f#27v~%eT$hGxRU-n zW6_k_dwf2n^uB@d`}(8Ow@>KUmAvg*EVfBv4t|PwS_queSW_jP#OcoJbVYQH=7rcc z64lA-UTmiss(08Ff?Jr_)}QiLZKjO#!b3{C)t()>bxz~?#w)&MDCJbn!ifm5)?v~#EVW8~Oee}Ek4xpw)nrTrnN`;! zx~@nG_esdTp!<5?B9&3{o@~mv==!+DmF;mf5udP|9Ges3$Qtz5md4#WO*>cQnP#7AvZA0!%=_X^u(?*vOB$VwEofAmBe#^HD!c-xV33Y%U7 ztX3g~;&X^iZ6GIo7SQwi>vPvnOj>$B?S1Haea!VwJ@nO$qd_GpGUF$MeWv%qr57Y3 zGj!W0W~bXEZa>m-AsVF#t9R{CgSt$Ij7niVNoB&;Pz*em47Hb6N6R$iNVtVRk#BR! z-3Fi}hSX8gELgRcyj5FicD@a7BaKVSy=8T^x%Q!rL4UA&*B_mAE#5H;HK`OA6>0M$ z+AzJ(7-yXdM$yu}%@YBZ3hZO!UJUAYTXNF8P1u2s@E7Cf+34qYzTUPRO8K-(`1YJs z=zP|@8hTi!S6tfG6C=U7nr_y$Zi1Xf0a+c(KHQxF96y)A4~O$o1`eJq>TTF`8BvX+ zw{7(NW_5|Yb*U|`C$C*FF@l*fcioYd>sIUSVsBp1QdTphnt4}_?&(=<@}wJ%c>U2m zz0(rWm8doHYRP9Tb}t^M4^CfAR`0~!;judTPa5zYv97maA%+mRtg$& zu5ybxS#zGVp0UOJ!3MJWCbnAOenW)Q;bWL@xDRfMl$0*Th^N6?!y6ycJ=kQWdCy+7 zxpSFvqF=#FM;p`XV!Ru7 zr>O5v(BQG7@-L;Kp}IwxUP zh^1{=JDsdPrP94L=(gw9OY(im_=;u5KyB?8q6p>+Tog*@-J@7H-b)CnwVFE8DS9h?cyz+!%V*wBQ}Ny>;W#!sL*u4)r$p%^ws8ntAM=;C~3&B!;O$C5uePNzHNeJ>&T zFi$-h-SkM|jvtGq37sTq?88n`AJ!l1?hLMg;3FLW1Xkl#BY2}Sv9C_Hg z@}{OLUu!6Q_40c=`Vbw@7~^$kz?eki8x^qeX>3je><8&4W2inF*=}=(N^peTtl_^*8b`L%xzQg4+5-{tZ&v)uKLA2_!`3ht%QI2#&P2$`*V>u}9F>lO5}6AaN2 z%pP0p4R_Lx8qz{xkL7luwSDx^@U7i}fV-e1__0d0q*d#D)y;!22W**hYS?gZ1JfCI zGWvid4@ZAwQ)v=peBD%P>^lfpmtj#?eUd;Ze&Fe1wF%L_44C%*rR=o4#Vjg#0106=@}Hgt<}Fn@7cm`IEuzAQiQq2t-Md&7BX4<4!mI zUgIHdlpQlCvX0e+8t3HHsC+1=qi`&Wmd>O}%)YT=2Sm4bD5~Y_ehb=~bVKA^(qlF3 z{-b$L7ZuJVP47YNz^@%4norh6AZ`QE*QK4m*CCJ3KZQlT&e9?@R>}8UnuV9T^xi;9 zWO|>#l^(<*kLsii8E4PRl^lXJlX_V6{tTtwY%e%_PUR5W=>1% zcDSQ+U4(na)R4VRV+2Ox5qH=r>-~_E;QL;9gf-PsaFdmvK`sAifHy0S`^HtKy1}wB zL&w$b?H)7&u=uc}&z=TKV1{3zi?Dv&EB>Rsymz=Mf88?3@r9)dE-NtPy+ z$iMP8cu(gl%D@eA$2yCHcG!le#q;SYC1-p4#=TaVrb7^QeCrNCSX zg^&QLe#W;$)UsKLNe8BT??i4~J-nu38_+C8V6O?FSTRY=y z0i@9JfOO3tII9z*r05L&Ge;X2Ra`a^>HW!mi-#Kp#&yjCr&Qvpt`GJ$###sEYY+*V z7fvxetJ%H^(N-_jWD#!EA=8q`~D)|L0xHkPA?gETwrv0 z|M|50#lZQ?EnKI4!WrAW!(>Xo?QOLk#AF3PC^l_XJ0C;O ze&}8YHq5rv3m_J-cIcRJE;|u#@vd(^kAcj}M$81lxV*`*Lwfiiz`0(IM+h-4e>}yX zC)SsT=(3{2*Lo8%(svmc66akh_5J&Z$SP+CUm+`gWZKdluaqhboZR8=lgG2Dcr}Q+ z9`Wqebb1XC4r`%#AN%#UOFvM?myKY7&nYrF4%$vAY>V{Od8EMErHEAAmd43GvUDJZ z6o6xd+3n_aGg2R#BMbzI&vHF zrMaM>qVn$XI!{*T;y&y|Z=9CPAnjcWe_ni^VZCqf^tu8U~);Ap?5(o8yh6ggW9MS^PhV?&d8!;KZYS!AbJ8hNA zQi;4C7y9{Ai9l1nC_SZ>ps`o5YpQ+v6xjIuRpC^*!Z}>IHW>(U?iG2v9hFcZRG;#C z3}zR>t+P>GO=t^=m^hN$aUJ>7U`cj;UPw#qUV<^rbuD?-x2l}LY>V0M8%D_a%l&x) zRQ*z;oxz1{f&zKfe&=4W9&Kq6+vRiXb=tk@<7%rOL^S5$jNa6o#j9O z=~Ds(1bY0facPz+SPugE2 Date: Tue, 17 Sep 2024 21:56:52 -0400 Subject: [PATCH 08/12] chart update --- img/ComparsionChart.png | Bin 22277 -> 25770 bytes 1 file changed, 0 insertions(+), 0 deletions(-) diff --git a/img/ComparsionChart.png b/img/ComparsionChart.png index a9bf25551382d89434072d8ed79cfe62b1bfaae8..07676cd364a7c3a723e19326adff8f442e00f4ad 100644 GIT binary patch literal 25770 zcmeFZ2T+r1*Def#B3l6!6_sj(tstnBNazA0(t9rogdR!=y{M=N2q;J=f*=B+389Cg z(xe3uLl2P_AP`!lxATO6`#t}>-|MQ)1{`2jby&do8dG4~-y4SkawN`ehrn(Z- zUl;zOp`l??zOSH7LvsX9Lqj)mlnxj<=Q?m0_=6Uvt#p^B0K+j2{Dt0DPEC%6rYP#z z)zTV#6{{H@ffdSyZ|N74v%cnP7A_QF{Z^7Q&_K231mA!ZGo~o*< z=F4<#ZEZtCLn9+2b8~YmD=TYjYe;;)y}iAIg98)_b#--x!C;=Ap5ETx{{H@ffq_q- zJ`D~Ie);lcL_|bnWMp)7^xL;@6A}`Vl9CVzL~3ekdU|?hW@c7aR&H)?US3{NQBhf0 zS$TOm8jY^0srmf*^OsgkeSLjXQ&USzOKW#;TU%RaXXp3t-?3OMWzVAoVE==IgTH?L z8W|ZGotYmS8=IJzn4FxPULX^R#M#+d5{WcFKfkcBKqixymzUSq*Ecpcc6WC_G-hZ`q>M_>h5U z=4fRhzgv#Usk^7PvrY1QQg?piclc;%0@hr12UB-6K=d>`CA%{kP#PLp*WFBPA}!4c zE-Wq0?+_vl&4xUnj-?@==G*c}9zj-Y$d~*@OKGOe7 zPU&g3I>%h`^B08Um(O!1c7gBwV6P>q)4OR6i%$Qwk(S(?T_P@7ky2$OdtaB2xeW<( ze!)hlc6>}+JFz0?qO0W0r7@(+!+~TCvB8Rq)?2}bMP=dsM>o^Tb@|%LhS8O>H}}4? zI$PYIC_aB74UF1+f5(GQ|5nhGPb|KU@Z#AY9Y^lI2fHP_`ip@Og-4u#9BcFVyW6t| zf|kG*6<+3YXKT0MnGJ?=W#kvZZw=tPx{bsvof)2JA*T8aD&`cX9ZK%GtY~v`V;MKo zI;HNhC4;XH5^6Y~DZ;xtp$jv}w-@*-*g1~q?Wtb%dnQ){Svv| z?+Tq%ud5+$eSOEnuFR?h=SAYZQT_P*;KPxP&%#dcjFWK

{STFN{AOFjw!3!UW+X}_S;?-HrKyqfflP(?0@5#!2%j3kr(R|ltk)V^3XOOH2y<75A_EJ;&$Hge-yF>@CEDEG<`K{#4-DN{>hv&mbKnD&=uJK>?^S&r zOG7S1nf3es&_@(lX!3d2Gcho?fobrhR(ZES7P9SAO** zb<0VOp|8X0Chn}yq|)x@2o!X!UzzoFhs!L+4c{Wy9esBk8;U_P`;3=Vo;b3;^jV$yK$a<0Se)DpmqwP=8csoK0OH~b_#3aFq zV4FKM8b|clc?x;j>DR7mXFl=-QJxK+%W~$9x|anBHV$0PH$?cqsgY{$rO9>qlE8Qj zkQ~9vIzeT=zX&l8wG9-ma=9xWBwS}IC4um=?~Tdb3@zV^Ox#pl@BqwHzaUaDy8PbQ zW`CbxO8MOKe3z(1tvyR{C$OEew^R&g=KW`@Yi8-=Mw{JYz<}X9{Fq2-8 z+0l7Ue!Ufy3c|Mj-$~e*#Njjbj)*C-w6`vT9R^vOY$93Q9>?@+zIMdS4EA?anuQHW zj6o#1EQrdhUAz9z>+_VI&>b?p1cPE3*Dr@X9!w-!uHJ;j2RQ~=+LZoGF?S6&7U8~J zKg4TrT^L{R=r{{b+rLH4)$xi1VLK%UeRv(B(B%{FDOFjO>dilNtHXj`<0VnSc7pg) zxcs#wDC;-eJm5j)`w{b9Cx*sp6|#XTe#<#g&8ak?sr@vXL*!Ukjem=zXyL22Q*VuP zFYDLU5XO1f$LEi!wnsxWove~`G0_&98rNJ_6W(cGUVXIPlycqRhOwuzaPoB!nZEL> z{bF>@#368||Mw22R%Wta(4)ESe0MWB4J2^c-j{+2m*!4lM?P+~OU69vSNlX}30voV z1Qj*%z-O9i6kzyd_MP1mZR-w`AHrEcH@O~c z!lo|m-GXn_bg07_+f~SeqT4;EPNF>W^6n+Yz8Re}(dQJdEBn%3(0q~@&1cLUL+{6ItA3OyDghIYV0+a&28@bA)ipc-(Zvz_R+W|ll-olG zY|GRs9(9> z2IzML&@bzD0~>O<;vO z-9R%Em|`2lSFi550mtFXbLS+l0m4$!YgW(6d*2R$#P~`p#`jHxJu9-CN$4-L5bW&e zXt9V)s?i7s!P!iV5jN%J6Q?{`O#Nb0b>^^BjX2{fZ zJw$uP$#D3l&zu!aX#RUQ5V44z`o6Uhdqa!xmfb`J`gGf}cQm8oMkU3*%*b4B&A*g^ z0KbM*hrfZu*+rCZcN`np;Lgn~>O@wDMU1{;Gw8 zC=&IX?PpQ1pK9|mzet_OCt}}?r*jGz0RovpsoAVN`%8O@!5kT)UR6Xf8s{`fAEe)G z0Eiz6s|gZEG{T4qd@d^x|Nb-8eH{68_m$>@86CP~qThKi~2Bq?rmp z&?S1dxmCM~iK9(T#mSZgkUr#k+%Uz@G?KYsRQV>YCdbL1o0YwH~5sBrYyJY>QVQC$ds5|5)$V*lX}yl_C=aQUsD~>m8Ze8dwHQ zJ)f+iGXG)4yG}jA%Mp*jxUBFQY3nVErBbnUVotYWAMa{a+FG*Hl5Ge#CqN5TNw@EwAJ6 zPGdvWKfw`)XRhAxu(vDkOzvP?W)dRKRrjfubeCDk!9pRJ>87~E;)Rtmvqxg0-9zF7 zJ@BePImA40z1Nrd^MuAsE{%%aB3TfQwGR|le*hCsoTMCCpGZ--2zQ1!xXv@64SIKG z>j!cVyS4p`jJvByez#De9j3bYSZ6%{^xV+aSo+VW(;HGE_^-E~!G^N3iRK&7#HtysLB@6yHca?&Z!V`@8Y zukXSC<&cIJO`VDjc}Q$8J(o5od;W&X?m+Fhii?&!`BqW}*|E1A@KbJsBZ99&q#4}(`haQrbPEG+I8LQ6pLd->f6s)8 zPJo-j6j(eAIPbC=;L7*5vXE%P?kdw3XVfZj&rMT}DG)HGr{cx~ z=@}8yYswW2X@iIT}JNd zwG_0{5gf-|RJ|!k0)Vvi;z3ZLtjmwf3mh1n{rmd{B>#uvk2oVueIe_!y4bova;-*N z;5{H)dIw<8QZhfDWqRfHjmXPRWd?eIUyCg?o6YBMarPE|DtZOtZpxXJu!IJYYKxV8 zu0ijTDT1g)=577u^02P5vX2i4>ci<(1S=6Gz%X0@#D@QM{UASl03- z-nqJ30G&BK%FUbYb2%*J{!sl!6(xS+G^#`JRLA-daI6x7@v`NRHwK%Jnyliv+rgov@}L$WJ3>{F%}2={E@2q1&p}|UBjnw6 z6-)Fb&0$15VdJ{YC!Z&T&QAuKJFSowlS-v_Cs0|nfh~~q-r_%R!`0sTcaH{o{>NVC zI6t!Q_!5YHy+Oe`<%eyZ2vCD^&!34C&$FFR6ct`CC)Z7yz>1d$58hWQbjPruYaq1( z^<|{=iJEmSEaj9WT_3RU~LY6)Y1zH*OkEccDsN0Ak-VoAMbr?&$f%pr4Y zCWEmz99Ax~U&2TK*glrjeGsDTL1h=()^dELpnVcuTqkT@-Nm^EH|M_22Iy^lt~@3Z z^a3rHT%XbPq|z>qWMF9Y#Nd!531YH*Pq;jCxN3btGO6%>^S%q$we z_%ZQC_)8OGaPTp|&?N?p+2>cJ?+N2|u(`CV29VBB-;xF5D(E8ql(E+C!1Mum;WtI~ zHT z@Z-{v`?pR5K1{E-efOyoOn2~_OUqt3-EQ#vwZEr4Fsgq$$0Yd8^Wm>9*05RLvbdV` zLfjAYS%hfZwMt588w5gI=GaQo`qI4FPK3xHS*0^M=6$t?U~c#Fu_XHgWr={w2=R+C z<(mY>itv|1Uh9rabSUQjJazIRaNh2#R;B@m;dKa--Q*&&+lTelrdDl^Yi>UW_8`ya z^4P;vwgjlZ_yx0w!ML-z>`_9$AH;dgpx&y4O|Ln4NSk}7#PSN&S;pe|tJ!DRW^{*B zeJeH>m;3KbA?kh>C%ekjf?d|%RhS;=T+y}J@{BHn*7aOEo@>CL>0sIPZ02+Y_+tev zPK8h+O=vp!x(xSEW98mvGr#^D8T_{vCy&-YAAYcSix2SD zrf~#5{lA3o)x0l5?0vNFDW_nHpr9E8js9s^B} z!4>SA=+iC)4x9du>}7ATmng5PHEMCJM_`XD_O@Agk9Zioa2ll**U0B8` z)0WUxsHBEIyzC=oDPcF=$-8%9Gflj5Y*0xh12}k9l!it@;2rU_(v=Dlk~tSl^9eJW zquzBYAjVb;ZiuCQTzXamKJ*CYtM-u zvuX>!yx;%_{;cF9whNINE|pa{(NN0XJB`GAi+Gal@4rguFzSsz6&d6MTbT=%&|03x zp$C5VIQ|Hg5#!t8WyNn3)lZ+@;oAOlTpz1+CdZY!~fwdGFe)2_wf$A33yA#$F5 zGz=AgQ&#a1Q9u+}KSlrRn!jtfZ_zw2*54hg8RXUNUguyF)=*;J9R&?rTk-;ttR}A4 zd39F>oj4J42Pr(*Qv9ABFLq{@R10+t;tP;H!Kt^WWK+=S$qoPrPKs*lqy%xg7l1P(hm$mj_vW z)Q#&O>YY(7M>!P^Y59ieaVjXIUJLySKcqV~cc}5N?|3n>S~vN{>9&XnuccJ9_l)TF z6~`N83aEae+dV;@8(XpLrol2Rb|9G_ zeHbiViZB9W@C%*FGG}k>t}aU@?j76es&!q9!nQ&5((^>iZ!9-+6*&R<-4npx+EkHP z2a%>2yjtfaUzfg@2hl%TGP`;V+wm|G%4D+8RGU|ITv`v=?W^Yg@sgr;n$R-d#WCOw zoTJ9DndQ*vNi{fR9)3Bw??#=8EnnMG{!d6FUl5w%WsAWC`Whwac+vdzv<4h_Zp`1A+4rd8oB3<5=jjEP4Sqpxgpvn1*?EklM8Uf!GuNP}jH^U&A zZ?Nw0EhL^@sQtRC7DBQ2w;SGE*LYMqA;a&?>Eym0!QA^*tb^&mv7QrYQkBYGu*Qq6 zOq|M120}KtaV!X5_e1{3II6A%PkSAST1X1?3^Dav1Y3OR(g+Q-T8Z9KjMBCgET3Ya z)p`&9JZbbE%+79+k|-L*l62J1;|VGvqOil#9xiPMGq3fviA0Qf>V0DPvHnB7n|`}` zfq~YI^40p^ZL3Rar;P>-Rs?l|V%G5l6SQ=|MgA+`wQa>k46yMP#|<8`#lW`1Z>uz@ z#dDTMH))1CyZqrB8*;MN*~|-x)(`8fU7og%TF7kWm^4|d@v`iF6|I*bSNo7O$u@VB zeze)ZJ@QMrjQ?%4>TVgZ<2S(Cw!Zy+_{yzOTKTiPvF9r4d$($eFBasf4~kBQsaEEw zuAXzPkDjaNHFMaOlXar3RO`u&(@k}VxcPx`#Y*G_RB_X7B{>hZ^&9i1x;k)Gz5IQt zSq!hDzBPxv@$ebg-o9;Guf!}S z?~EREEP%<;E5BQNIHLBDH@NRrtn%kek-Nj02@61O6|f~-7Hd)hgp^+ zCPISP;W9rQJ0sKB>~FmP$wwk{Cs{&0xwFS#>+O9O7qb0ITW%N|JA1G}p+~ndaI})8 zZ<@t~44-I+Sbo{t9f7o_Z7he<<>8Mzug;va7|5(CNT?`>OTP=46my*(ksy~lA}y-p z_*Z%DYCgsT3Z7hFQ&XQR=*DK8tL9&kGh7LNd=4-^nEl1}FODLASybGeY8JtP!{!>Q zT^=pvo+c<(I=3S!yZ-$HKR|-Yv>A~gD3@{H+42ipSA1H62+`d14UYjCsj&pEETKUc zlKd?WnX$LXbge5-)~(HTIN}m|h8pwr5oL})0*;G`nb!o8K`@>}emzkGsfFfsw?Z~N zls=emP&a}~%&b{Vzf<-&FYoQGTMW5H1QX4CK&Jsef|>6KkqezS)lgOeTK)EEWCq$> zY2e2UGKtHtBj#HpjO+X>i!5rBcQJXNoT~d=`t$RQfi;sl0?Xrv*OuTmh11;5UbT|^ zL19yX?H$i%{?%AGpM=lUm&vp>%%>;!UXgBTs9d!V6ivP1Y`o4utMTN&ES&OR{LMqk<8hBx(GJOi&b42niXJI9}?lY0Gv=w09)C62GLPp zueNTRm3ckJjMOo==&G{QZ&3oK2g-+ZdAxDNl>4gc4#8D)iT)o%lisyV{Zyw(iBd3?yuV`xpJt!(JT%fo=USe4G99irp2$m9`b!2@2rXXH|kXok|+bzhG3IangM1zm1F`E zW2|zuvXy8X#y~(BBl=?JYA58=Ujj`G`m7r8Fiv{tWG?Ee&K!$PRS}%TOV>rwIRF&h zs%Zmava~o8-ND~~9VfZL`rIxQ0X`IX7RfG%jLaK-<-r|1opZe0Kfp-Pu8m2a^Zi78*M0DH#*C1skeGAWppb7B!p6^& z>Hccma9O{PP<=kbyeKWumhc$cpp97MX%k}+5@G?_?B9^dG z(5DVHAA4YB>Ltdto0I(s0MS_xv@5YxN!|Vjt)|` zn~j_?;vjXvNdf*UbZz^X(KiVnb?j=P#q{$!IIBrv z&T>jVE22m>9ZW;eL?}$%D-z~R>0pypzDO|s9h|bvjyltkSGwiITV}2^BR(hu4fp}k z;-rCPt6vC=QQX4MF1Xy@>R{%t9q3yde{nc*C6+xf8XN1Cjc$^>+bJrJN+ zborPp;%7Lkn3qb~r|%tBmLHDqS7m>wUSdYwkpJ7|lvH@F1cEBmA3Q2+kGAji z>JNQd$f-h?kY|!XL0b|};IflJ8^qkXA-mKh52|UXyvo{RpFcfxiszDygX{9>QriEqZ1$ z52^mv@!V7bk;YCA9-Wr&D&O_vuWf}%HruiO*|9{f`Hwz^>NVAfve~3on7%Gwpm*is z-9Tn6?e{R^m5285VE|i`A}7|-dA~(gY&R8oD>YsSR(@<1S1AoGTnhM=&HVVRE+6AS z7>4vnuWDFlP*^Q@F=9}*CqNAgJL|yOdBw$3^>P* z7#_hkH@biLiA@QEnmm$58seJH54hSg47?m_5x-dsSkR4Z^;2;{Whi1B@yWoB;N(t$ zs&ccaD8?Sk*P0`=+@=@{{T34hc4;x*i^F`3sn(2mABB`8q3gCiEQ`8cEQbDS9k2@# zp&JnrOytzR%?zm51=;A9yn%ps-RV6F4A&_`Y|Vqfaq;?wRuvv;lR z*YwUV&iUoQGvfV(C(1!CvmqWHo?4e$zI;9|9eCrG0~c2o)$ofNhUsaHDOYM$JsHc4 z%WCH|-{@|u7J{s>5Y%>~@Gal{<(C;+I@>D8j_=trn`yP|o@l9eH)wgI6xzG-HK1Jw zRkWCUey0G>w>`1yGCVu8?xfOw1ufdXfl-B40-i#!!W*z-z#_C4bGa-!yzS+!YKJYy z6azUpB3_d71Mp3V>y<9-h&o=c;t!J^_1a@?4|4W2eV*H+J7En2zc=er?M-@t>Q9`( z^vYB+u2?I+i>P!g=Hmw}LMa{0xV|tqJ8Ra@vG|1)0og7fhX?1u6YF>Z_;c+482tH9 z1zkTDWp@0mcjD0xL#dvLvq&l~M=F78biqU*i~+z<9h|Lh_-E*h9(=BO%nwB0a}2A@ zfMw)=M^x=B6D|r6gR?r^0}FvXSfAKwB0Mw^pVD(M9g!*z$^dkxhp0$Af|UcBk0_HB z=}wv(c!;sEfF@EAp~;4B*6w4PoIq;z0a){l3f#kt@Ifq_9dh1fo#&A*AEz!~9uTl) zz7^9+ysbS>phG3wK@wL$=2jwAMHW0t2fsjG!)}V6(>(z zM#eBBS@o}A&+87ii?I5BG>U&xs~fxqJG0Q)J4>ccv%jwD!&d zAPvL2!k%OOu_@m3P2%G9i?chvB9a`q3T0VU>Z$&frb~%A6*0_g{i5#%hEXAyx={d; zT|PC`8VZ3ALlOf6e_k-H%8?bJe!Hlfb!RK7wxv5SlIw-Dt8akJ`I)Ty40$8~&+nbz z8#(rZgcZjWbFRL?RBXvkyYH@pzK7RU999t6nTE1)36?eAUSc52P+*rgQfIcTM;~$B zXwY8VTMjKV?nURjU|at9Wuwdx>1%!Cis8K-%bk(9GXC&|07O|0BeLW{ZY@e4{+*Q~ zATp1mF(=Nbst;zI3nMk|n@G})M3qSd zwTm)pi7T^Wmup!T=&8a0s8!t_m!Y07%J7ziU$7lLNTZf(~pU!|z0e6#Z|y z_v=MHf@Wq=WG=|hsEqVNR^VU;B~UI?`$d^jnK}8iZj!(t%25_Pc;-Hf`X{2GldJ{x zwIW@ms5_s0{C#|?uLZdIeVX3+3GFRg^&RirCNHk3_IRuzIUD|1W&6$#^&K^dUf-_) zlV9|F7CN0e1A0YIuuxAVf=#T!m#Xyx)Ra<>W7y>a9CziM6RKC_k^=mP9Q z=YwnhB4C$5Jo|~m)X!D1v^{Kl7b<)&dQdKYZ*QpO)5>;1GPS*Egw(g__JHjs9KSK>w3hp4Xu6+#AXo~RoVhEouw=#BM2;l#s8jtr z4K~rNYSt-9<0S?HE^D?iQP9!$-!A3qO0x+mpH`1E>qAvtla-{-dJ!q?=Mq_0mRdOc ztwI7!K&1+5mGRHL#_DCQkhmzd(Y^Tn6kuLdo5s>Fp|;8HL+ZIYPEXO!`H?>%D&^?C zBI_RX$Of}~z+Pjn`{=UNu4h|C8eV@-sl4vvoQjdBM2?A&(*cI?9kNRrTaeqamxf#2 zI6+BMrzo3o7DCX_4^JB)l=nI3=YxA2LBA#Qz(L>UkNp1C38g(59#k>-JS3b{^lIzN zR{afU1E5A1whAxdk07d+g#`@lm3)Cv!DX$Mn<1gX=BMsxaKXYmAi<^?y7q~v$|ozd zLp*z;Xe}OX=3UXp)%K0Mea=O}@#IIa0fW4hh@9T!NjHEjNe2UD34WS%yUlY6Kl>0% z>-%+>e%jLI>8HE8lQrww-(O&esuRZ>m3)7{1Vg9#=0Vv;OSC*`wg=KL!dLhF4EWuN zSjzSO&Hl*?_Z^S|kn2*=o-BN&^e>6}ZLN-{LPD=R_XggU{dSAxJ79 z^VnLGv8oSvzW_>otI7QE(^X3KngYK7L;SgN+_1kOZ#F?b&_NFyAyRaGoXNGQri&Hfiqd)m&KJaq40=20bh9x)7^6C#b{OJMnY*~~qM zsQdi92{!5AOf|5omW^39dJPW&OKbt%;odZ7!KD zv&%dpR3Gw~2o%t8A$3J`d7=j+?NBjP0Eqq_BX!^pXLs3068He=+ME#G%&#sAp!iJ=6mPe)h?MQqDuSD1|b(mlAfcJ z_&l)h7Yi|Ckvoiepabzub;;mPc#BbI-N_|?qs$sXUGi4|LWyge)$orxeSQOTFDo9* z&8W27K$buFTR+aTOv7`bn(A)u>W0kz0&JcFh7V^)awS=EDK*Ibpsw?8JTfhEjCr;H zbC=AUeW|^!VSH>$MNI$r?SJ#(pP`8&@ah~Mnu|DQC`@%|QP}g&b%=_QyJK|HvER6e` z;`%)O+%7qaJH*fX$tO{dBb_Upj7xfEccoA%%{x-4negH2E-?P zB#p6qaq&n|_lC!w?9PqhG)Vfo^#H?;fim-R2xYn z(sUl*oAZAO;_m9(-iCk5oLBZQ=Ldi~UHiOz zk8|N2bM*e*)8)J+%^vY0B%@0d&lhIn|bjWweOY|YBBm^ zu9tuHr|n;Cj8B|d<;4Y4>ZJ!%U4{#DC-Hf@i-magPT|)_T?JcQ59% zQ;M#8ZXyRWd!I!^v;6nVJKe&twK5Xwu*oDevV&GJr1>qwNva0}&Ylsu2Z)lgtf}DmFsAa&Ls+Aov00t?*lP)(ghGjSZ9(wgQ>oAjuvxrY zegDc<)vLHX!`({5haq?Nqou$u*5_xJ3_~nX(;=Wl&quXk`HXt|Ve&ixSl3MbWO`zFAo zs&;hVo$<%KQ(q${hd5L`&SDO~qatnKm=LeQ{+;^b4D7YY`)O$A5>LhQ`cSi4ZuUj0 zb^y88qkyM^X>QZmnBfQsMulewX7!N?2KMsoAyzz8g3pDl(6~tCa_^F5c8~Zjvh7g? zxPPdE=6$##MSoUb35QVeIl}K^wFz_{;9XME0=y*2E`ws5-cmgUQLmSL6rB5PhRA=gP#^YpTQ%>Q?835u563X7c`P?-fd}6l5~Hy=7UvbNgn<*3XX>%SqFo zaim01p6#9Q&IW1hiB*jwGAF6`E3TP1gpNJEp4AvF^r*YfiOa&A%jH4f5SLSxOsPp@ zCFi$=Urv&Ne!A-Pqx*Jh$ca1t(ZWB&D4~3IMg|^a&$s*TnWNOmAy}ym7MV z@W}t0OBy1QyQYSio=m3p7I_$ifrL6?dc-sswD`Up?{-6n>+3F2bLt$^Ij3fO3isQ92! z3xOt{Ede<2nfi1v{!g?T&5nf1@k?*{T;;3s)Rri4ilDnh^-q8U-q7#JXHTX#t8fc% z;c2@EJ>Kw?8fSnUx7-1W&V?n}m>`rK>E6*$oifn9iq4!)mJH9$aJ(B2J^l|%J_P(Y z2P!tXnmf0U8bnvK7Ea-uQ~%+dcgHBabr!f$A5}oK^YG+eaAx>lW!WQ}U`{_)NfEFv zL%?2CWGO!Abu9%qhIly?FtEJBIDu(>7Y!(JukIg;Tzc*PEoxgUlG@g)7ta#O{GDxu z_dbIk>$5XT@Q*-yhVUIeBQe+d3sSsaq`Tx&s4{lbwJnRWkq4!a;_wcdqR~?v?msmY z_A#9@oRT@8QLMwJP4v-iF!eM80@x>-mEj&NhgLJeRq8NG=IjVc;3_LW`jVb1HlRBJ zQu;)EJZpt-_uP~^CB~#p*Q%24JYZhf#GNl z=C_js>!7U;=zZh!lo3qClLYFMy}G8D5=mgTby*gob3mK|h0D2)FwB=HEtm#1r*qS) zf5^v2IhfGj9c9P)eJ2urk4BU4}OYp8=F~KaiH}(X!V7c7#tD8Ei<+-Y?}~(v~2YH>aDJ zR5DgKg8n=J%9x0s&N$ijn>+6w z{qxn#Ud}y zNS(=B;JJMOJB0n_tBTEP<8^ST2qKTefdYh+Bb3tDK4n)49{U|=y8KI{{?J3JEj&Up zfBKJWDu0N$|7)GlCGvE$OKo4D_wt69S5*#t=j*ZUH6nHuY8S9Jg3V`Ql7j=->^`?~Ojp$i#YNiCd))gC9w=K50HU|6{gYAZq9ir~n-MYahT_pLQI&Zu)aH4{O!Mz5HV=%-hAM673uer;U&vuqWj~=vx@Ld{vU@u*hU~%3afe(CdQMak( zl~JIlr0zfYB0_x4f6p!jkgRWn8SfYAkp?3oT!wA}lalX-5B_jpUgP-lLNCz%H}t() ziTs?MU_9w%Zs!KCyugrxbU{BQn+VDZ(%P zFo6dK6csHFM2s-(OIhJVA#rQN-D)hisYk+}_dxeAMIQYE8jA2&d7uf72)xhx9%6O( z)M*KmdTJNISuC*I^pJqQKp#1eLsU4eg?K3<I5t=aHGOm}S!v%o{;TJs#<|3ySNs-( z#~Y}`T&Xm+QtE>Bqul|-kw?^&&^Vb`E!v^Q2mR~B$vsQ#snVT%;D>^jWbD2I`rWCO z;5^YBj0b<)kt$~r)T9_@#dm0tYJmBac6CuUh=(xHK=mijG>#$80!3&hjVrROpd*i{ ze#>_gLuBOBh-*OMr z(&WV~XJeF0!NsFGbN85N4EbiX5amX3xXMK*;kHQJ37SvlJpg6cps}gQfa;HU)rmVt z9jXe9nFr@vlFWxg<6ChfOu|GVr6zk3Ap2 zykZg=i|e(7W;HsRykvL|Mi*w~j7QtkJcgqLfpqB}9Xp~vG?|9RP+~?0(Ler3zyB^Uu?zfoH}b9$*gUDEx3`x5HcXg=Cv#x64_% zwpETPO_Nox{Yt9q6q}{C#0IMqJ64L{OVmWSbtST&k0Ez`arFjA^IxBGzBwETG>BZx zYKGcF>Tql`#sDQu{dRYN0P%#b;w?C zcxYy++Dd%-#XX-5|KB~#Jv3M&mflSxsG*?$&}q>-?eW*y+>upN=RxXYFs2Ev{B4rg zP(9kn^EUd8ux?CI!#0nEr@Wh2@ijKkiwTXc!Po5O=ultzB+%@) zed9AX5ytj)u0*t=gf&OquBw4s*Mj6ZRBvmZyXB2r$+%AlpWsM;|gT$?2Wflwk z)_SV5^PM$}9#zGX(C;OaE+9shH~UGAi*d)z6n=}%66J$LZ+E_AmC(C)3<=Xqne6`| z^XxLJrM37>?OU1rW9U(>^O5BYr4q#-VGxOFqlxpEg!sew8Weg9%P`HWF^*(|Ja%iQ zn##$I6RD{VOrVaolt%2EQlgyA_B97xTc&Rt-eT>kTJcrq+&eALPiO?0ft?c#s&)B9 z(C;4NZ{FG(uAhvZ>2Zc7ufy@AZ#g;=af7ex^nRQ34Q3`BKMN2erP9YTWe0XoTp0)< z%i|reZ{p9Op?3OrQQiJbJzHncrBkgPhzB2M`QJnmzjZ^*%ndXVYq(Y&w{rj04Sn;a zD!zQliLmFdq<7oJ#E`5>0*m2I?B??ijd+>Da^no@g59O;)A>koNe|!Kd^;Y#w?x#a<*3ZtgelItMe_(0MxH6O0D1YU5RBuf|HqO*Jrq*px3(;cf zyti(PaKVa4G@o;u%HZ_JL83hTZcJu(k-!x0hn&d#NQ=B!wrINVC^U_cRwlt;8 zMM*a|WBv+DyDUTqi3x~vMvn+%cS3liKCRk0PnqO{mj*84Ux~#o{~UP7c)L?dwKUm} zkFVqwkR1<_3MzTs7nf|oMBnp2Q471i0v|bsyegey=7Pod;c{i<(t4iLwv|xQlz~a6IOGtsh7~OR5ikywx#beQ?is0(yMAMe{QC9tAefJu4fu}N6?;9Wv$Ukx!aD)VyU zM%RyMf(*WqKTFc_#_r5GBWv?+vSbFBV_gbNpGZfW0qVQrSS&+}4SlEoVP{~S$Kr8I ziu*;__`7z%Ch{@5=5No=uN%ybY;OEeY0)IS)9V!-FS;#yb2IcDL!+R&d*k{g)i{mi zu=WC_xeQ5Quh(Dugo$(tWrY*Jk<$asIy>Z%msQAHb2ZmIP4R9~DmDbTG)y>%AJhKx zT(%i1BE8yyY-y^`^_JCVQQObx7Uu-_L=N6`)gxE`9#-8&tbpHGm(RC&w0io5+R|Ho zwkz(^$nInLk(DOot7<+hj<3q11e;tY@cw4sFj&!_kIdxb^dKxoIQx}_*-=Av7xdv5 zc>`yYO_Nw`6!XEw9H`zKyMlR`i0U=Z`pKLLmM;a*XTKSLgWLQq-1_O24bk$>?albg z=4@%$=5O~8{PIw}`?jq?=#Y@xozM?TsHQkDJuB`Wg0s%Sil>j80UBbaccr{*UIa z{h#Umk2~ocbtLtml*}ceTvBe+42MKf(%f2`$X#wH8o8&A6d|{k%giO34QuAMEh;5* zoiLP*4mLJJretG$Kj**r{{DWv->=8({dm1!&)56$dOlyGhg$Ma;@sm4SITP=UQy%> zgx?Rzw|01wzcE7rTU85-P+FqcTPPD0Bi^uO-FmlZt=H30@56C@u~Fx;drcbY*jPr9 z7#tg{svftSm2VTczcE~LPHa=UB?qVQXB_&>Q$!UAMyUMt3lc6>Z3>PVsT1>N9s|;C zDCk)N`L2WZnu86>mdvD}tJ@|k%HO1GU06OCXV|A<5fvTqzJM4?pLzerLQ~dYTp51w zj*|gAz#ztz?&*ziBlvl|?#KNXoYcB(3>vaor3`4G?Nyw!fZMu+)+OJ4bz{@) z1at3zhlA2ZY1&`9EB}_B4x3!74 zGqvCVgknYYGoDYnkl@R0F0z@WrwN+YEzs3{mNk(0A`19}u&6QFb$MQM37jJw<|`5n z?V2It*u@tkStX&a!m%}-2!z_r{;PBb0nq+(vk z3*A-P-#jVQ$L$wrytl8)Ph5OlO2cKB-l^sb=KGY7j+g00QTJs*&-TvJEYB~c>|rxM zTWcFv=weR|Ohm~~2N!soV`~91&!Cg4`TQ&Vy1^RSEpc6=CL#Goa&M*SD;%I{+p~tR zf=tF}{=6e%0@Y`e@2yG;#RXLrX2J&7ZhIYMK{_J%Vu|=Uq{aQL%LwL~unCuTO#K6J z6D8k zFEAba&1*82Dif9ErAU`F_UMmy(xC{HD`5$%+LagDs!y-u(q%Hb3qkO-&?udz>J?mmzlxxbGu zt5r=V$Qnx`vOfhXCq8#X-6fIS>+UZc1K4vzzP@LNw1$h8Mck=Tbn}F0AD!<=nS18` zYDU5G)1_@8wH*?*%_l^%dm0jdJ8qq1X+>3~`{Ap8ePaPYfy;s=hoJmX)krm_oB|X5 zxX`N^5!q1oma5YVP1BWc>(`%}$~ww?g4OT6?frtFXWgY^!zENPz6+hI`t|tpNm~Xh zcqU~!Z+v%`(8eP$L*DnWLS|}&ud&0UyM1PGSxk!m+-JSb~V zTzw>dRoFWh*aAhAVP#L>3$jnJx*v=Ewr3z?Gc`b5QVf; z6ymPI;`ZW|3>9j+qCT`;!))+8N+rTWD;VfFqJvrgQV+~IW+84FB;TCa@T|{{lT#N6 zD;inD4)t)5*T#a)3FMI<`Xbn(uX$6ZJlv~RKXk6Y%TrmFt?p+$d=umaE6g-ucQptd zuisRf2QS0p-B8)Rn{e>)Q2++ZtfA#bB>xT};}TGF9gJBmf?lD6^A1!mv^7MfAEarv`_u#9&lNUXE*Cb zV$Zo;0637nNommW?sJ z4Mu=|CcGVr-VCqoX(jg+QCBBTYM1Vh0{lcu>1wCdJ8|{fLUbxuDg!Yu-bYHQ7wJ7* z=Pw+*MSjoQh*;l-J#@G1`x<6j4O#TTUu}kw-16=$j$GdP%aI{Am^;8{o|T!to3h2V zU_W~Gp+0QU{!xtYSBX&?FpqZ zD=5G5M3jG#8)Dr%q{fYO$?bX4PxhV2emx&Q+E#gYCC~sT5MGgEflnfYQJZnZA65+z za!B${WHyTpFAA|fROieM^1z;4w+Os~nya=KnP>1Ow=O*kNarM&$>NPmFnq_lW>^ckEgX;BzvjaRDEWK^b*kz!FFIP|VCrA7bUf-D* zkFNTaR48$RFl3tlZHxakh3&yk^~9|nU-1DW+h z!vaF6D=scUSb7!nqr(6)2vhB9&uDR?D$uNYEjO*ZQd2!CyXRY+n!L{O;T?XcgST^^ zdovx78`B-BSGNG&G_wkQsuQist(bt9VBBl>A3&dpe5XQ^PST%4U64FhxXG$#^!lRe3AuT+RpyDT6Fie7#W4pm~wjOM|YB1nI`It zADvlxr;giy;3)k5#`aU@-O!}xml*fo{PDe=vX+>2)58U%Zl-g7+Axin{(&;Fq2QIsN54Pfl(R*F=S5~L$lLO=w9%0@s0HX@yXf(Rt^ zPz;eEDxCxs1B4cp5(!BVB7u;E+!Z2w|L+_BckUVQj&a9*XB@{+aILxKH|JM=Uz=;i zni?B!+PHfo2n5=6`OigMGsCl>GLrlx@SipAXN=B( zK$VG-{F|biWu{jxF7)^J_x1HrD3pPL0pQ<%{j=Tifr?+es(*qyJn2N>6VsSH8yg!3 z2M1?oXScLcPft%@Utc&Je&@~|BoY}J85t836CWS{l`C}!djwrYj@*kIQ!qs3!PBz zZh`05f7MFVe>_b^9P&D?U9NmuXX{!1xSEhDTnlpiUC4~!bqXh=K_CxaSA%v+lUIQrN7F%|mUDn8MxR4K_hw8LKyU68Nr4>o|H}-6AGPA9 zfxq^BF*bZqWGqjpYLMkfq`Cz<8y0@K=oZFOE4ouAYgBJB9T*C>d>MyHi_cx-kq5c0 zYk|_y3#)~O{#InXzCQRa2I6!V&yiVWU~yO~8V&c0uH(3yfXExp%o|S^8K)9Q1Qw?^ zkz?OIr}w9La~vZ$YfMqhTQ}`^lIIvp4_K%X8l zhknb4e4H#xkhiieb-_e>!&JM8CThNz)egbs#rkR!oZfCQN1vNBo;sMjIV9!SG=DAp zVb?l6BXIH$Xx4rNj$hynV=Iv9)VRj>nJza<90#tzF==f=xe_ONFd=vg~Yano_$a2 z=5PD>k{revxoc_60i#?8Xt09hN8VmWsQ(=F$2#S74){SXz>gk#)I>_a8LOS$tXRdc zqdB+ws&i*+a!HDvq{#dc=KJ1jS2@>NPDCy}hf50UzFXRG@tCSqE)4f}>7ic5A9iGG zGZcA}yLue`?8?=$Cvst>NSABMIw!y3${-kGq5X*{wspi@j%cwz9r!cE-U}%z`>dH+ z`BEryWiYEijyLHQe54Ui`ku9R#xD2N;N-dK$cpJEvha;uh2ucPdgMlOEa?Nit9H;f zw6yMc_ZMPKqociM-^uvxIX2^nvLpuNdC1n(_;4SVHCOsX87|nmF3^(pid$oy9uMNd z{2T_uKUty(P}8#Y_3y_v!Xt^h_<=nfE1K6txzYm4<{YMHq3>}^y>(f?j-dhsT0vM&^h@gu#h)v^9<~}|UQ$JibCRrBl6~CZ)Gd$jpc_@01`KRYhD3S3~^%FtDTDW@p#{&xY+z)t>ao{>@L#uYID`8;5cU<4L zo&C{)jr6NCajQ|-$T*XR4(0Av&@oJh<-t_PwClz25EGI;4-}cth>I4`MYah z^C#F~0TffaeJl^q5*0vq8^29o+ht5_v~l^Qt10QbXnAt@VwazyoFJ|m&|cmfxhlsS zfc9<$wD$-7YNz|hmIeZ*;q73*VDnP=&3C7KFqUZB)i0VRHt(L|4W9opFj`~Em)*wQAh zTCk(zsC-F+K!Oy9@TxQ}_cc+SBVM~=jR@e$o%%D#d%u++PQty30E6klvn9Q__aY|V2L6-JKOMRFs0$;(9Wrt7<;{5 z?E}n~OSlCo%T{AQ#&Niu#&}XV$E;hvCeXs&*}AWJUBP4G%V&5;h0L|V#o7|SxlSK5 zR>xVi0B{@L^H3#y!Zv8(vc+LYB*q+eXwd5FEsSYh?1`BjL-D;1u8TFkeOo+&j7;?| zQCI`BUh8%l-!fSWhhX4c{CcX#Yd`0R0z1z-4E($FQ^_-F>2^groocVDycD*0!gJey z?vS!-C7HsfKKr6l@ZZPIQbnt3ZGf7|UvW~?JcUPi$TGH~-CcQ#xhYB`f1Iqi-N1?RcdlVktjEapO(VvY zYaOkf$(p|T5vyL-}kM_lEl#VGJ1!=b~Sob)_q2L z4?1Bulx)W_39!V)8LR#zUh6bQ5Qbj08~I4O;%o4qosu&r`)n10Y~7oO_w87mC~6)tY@VMQEARB; z(fT;&DOZpT`Si8(O=SZ2+j{?xm7KLKB>aK(s_q?J0{Fphr?*XY-12)ktN~*bRED0f ziW3}@oz^3JmORiyzH68^E2+6bKDi*#p$l;1HhUctTxz9IXZQhjq@Nsio7VQqFX0B` zy**jokCO+sjdp8duC95f8Z9fr3ip9v+24`d$gv3*5sZ)K#BOB$j<@zE&iJ^}JCZI3 z$`5uTz77z|G);uGsy|Q`=uxSP*&Y8AH3Y_4xHy5|fh@2evOK}c%1>p{hb#tVI91EE zRIwQfK_{z4tTyN&V8yJ~L_7PzRi_zNjI!A>dVk^VnXd++&mrIU&f{-(vH6uV{Ua&)MTjf=Nkt#SK6)V78Xs8p z#KqE-;avx=EK*pPGC21qeCcyTzI`!-b=pxV6QpD*;_iSXlF{!)XnfCD3k`Z`dAPBr z{N1r;W%;6n;nvTmvi=ktj$5tlG6Z0q9ywz~6BgNjOwbR}lnlDdY^Y6rC;gTn(Kq5r zwK=|U4KT~y}iG^Z?!n-gAr?xl|J^ap?! zq-PR00j&!CTdM|YNjc0XQdQc4JQ|Y~H<7SoX6`ym5i$iSd)9Cf3!Q9DXmFe62MJn& zoXgqJk2R$|^Aa6D0dc?70n903IwpqN9LO~*IW8A5d!OEKem=0wE)MtfHh?+R)AEv*&Fkh2e*g>YyHoL!)2NJmou%k)i#6@M-W8~#INK^eE-<`XElsJsPi}3 z#K_R(JZEs$ABf#x3D456QcL$w7yx!VNt6RH1Fz`)`mP_p9T07!Y{IL7rqdw z_V@yTI&1YB!heU47fIagj>KrKY_;{&$>5xW#A8`$e*(OOd`4P02{9iOs^W~KF4lm% zhK8H0{9!fZ%AWHDkod&NYBcdr&+nCiW^?NqFe8GVGfCeIikS(qP8Aik{NnTPWXkSR=uzpPzes}x__{a+FUI;<3n)BxCgj9Y4l#+;Q(c)~bn7~hQPXM2?>^3BpXF2eYmm>c34%u@P--za@$uF*Ul?S}itB}Z2#B0HEV`l#1- zIBC|5be6Z{hVY>PT}40kdec0SIIey>;CXOD=D_?%DZm68&SAe0j7kY{MAmkdvptQ! z{Nyk22AlxL*Merwo{DUckf#JAspIED<`lexdg(e&u*Saf9|K;RxI07qfyuduFdyrV zaY@g*adTU;{KD$-vmXvA?#R^XztK8KEob^x0lWE4=x^wxC2VS0Qqmhv6#QINh43B` z20lH#(g6Sbc||nYYXO&3TRBSD{<e%osc!(RtX}sA;KY|aEyOZm?>T_yL1Eba0>mLC|V!w>kx0eka80%4OqI-RlQvQ|c zd|w+0z5pck7N{Cqzi!{ua_KAh!w0iI&p)jLa)FEN_1cZ*-<`PGjjF?t0>sb1WQrg` zg`Wc}uP?Cuh+d(ehYYnJJDL0;1be9QDxJ%@OIxvNf_~4^Q60|Qt8v*py?vxH)}ij< zs?b_PSjw`!2{%E{^)wS_JC<&x+V?ta>p#Zt-vC3c#(oJd&UmHD+JU?)C@mRT;bViC zk5AF&`dY`5%N({bRJHRL+?2bSXNxgG-k;Z?6ZG_pHU0RnZ9T~x>ip7u&c+oF3H&lJ zb!YHur{Vi@m4u}H)+wBr7oXAC;WCy$_OZ<*Q*cqco9Q&AdvuXl&Ww+X z?u8V5aLA|Ra-3TodVut&o;z(l6R*i0uA#LKzkMB8h6Wt!z!bl~vM($mmIC}jv2QKU zRb|HcgSTKAkNU=}gXF60Q5?0l3a;O0H;-P21VkLq21U|vSZRElA)0>Tv(+s zQ&2owI$IZQhxiH(^pLgUPC2*Ezu&%+tL|ue(6re?;u;#gdan}jGY@h)uShEIJmTK@ zYYOLM>7dv=NX%C9s;w6gn-8H^AACgL`-g?ZrETOhPm84fI)aWoa!q0HwSOKw2zyG6 zxN8PF3_SW@yoR|_tv!sUk39tc1~S`bg6Iv#cBK!O3bt#Vc;0RT0%bq-K(p274cbWr zzU6jkLCf;B>uBtg?^nyq3n!rNdbWDnnwyZnfxf(WZ-&w{e}Zpugt8CNkn2IZN=X*( zTJ2TGbnILxF5?e}$J}9pogja+{l%nw%G((pnb7s)<0gD%@oLcFV^8Smo3g78vIDS% zGFcFifBpXA1f|PdHZY0W!9bCnq)B8s~H-<0^#lD57$w10R^u>06Pp z?^*MllrLnvJM*LVZszW_GtQtlv9bA(Zbi*pq_@mZMysBSdExx^)u7YwzZDyYTzt?? zpjQC23zc)|PgRn1LO7PuMhYNZ&l!tWX&wg7cu4|J_^k$caNm{tNW}R|+nOK}Mm_^I z?Pp3+^q=xZ{y*^$proyd@<5C^PR_PX>cGDe6SG{*-6vZ?_hhyiBE~9ePH0wrHQ#e_ z*ijp?X!83;P>iiSnUenDshoQ7OBH1Q!7PXbVa+Pg<{Diz)}XU03hdkMHM&k40s8X! zf*C4QW?Mx1n|i{qdFx9lQ1Oo+%n_s+f{DV_0$d-tNLYHxW1Ol*yH?S3;Y0m~9fl z=xP1qw&(gpg1h$7!kwq8M|sSq;}N7P#wTg!Lz9^xY0#JR&E_a32k!Y93a%%03cLop z^{GNjKhV9N&E`PCzc~NDXu*HQSbnuO?rmP831d2iv+T9+AJRel_)uZ6-4YZEJMCR0 zwfQ&ls)x@Zr`KCZ>=fNJRdC#iUj653#J!vU*WJuRla6*XRw@ zT^i0y6*Jb=4Js~ogHm{%r_YP-Xp*_QSR-{26l|K3gr4k_^@>o~ya~PfcbO{)tE%T3 z4TPLbVK&0THC(mfWTSmoSA^=kkJKt?6dm+V1P(IE_?Bmr9d_XdvQX!yK4ecm7Y&O0FqT zFdTa8;9>allytl@{^vQ7PuQuLKnDuOJ+A=+S5%+FAQ{YQW>BiILxF=S*VRxe=@L7qVc3t>1Ey1dyw6#%r)? z|Ij<2S7F^ecrxIR>>bckW76!pRRpnqFFb_msX1<-PR%Oa7Xxj$d6yEoH%VHovK5!f61#D z1XDtE-cddp8iDFXY*r?>AGXnMK!)IQGD~+ytjStar2A~~LW@p0myOEa6)jV_Pg}Ii zRC@ntLBqGKps`53@!&nLDeMZHp0;^8!R=Smah?%OnEOsb7%!0*L~9!y8C&Q=AsdJb`KGq4;$G0#2&)WSky z{e3{+OcnO7`R|+mHJM-x5&3~G;jBjL`elg!v}G_nV*qtoqzT55q69!K0O0?01`!Cv znQPF=Qt13+XnM(P*m(3C&*zZIXOMb8u_kWj>|#g9H|E_`>e(n6aR>m_<&l~6ux#!N z7kAhax?T_Z5eQ|a$d@I^AGiN2-j)hdv~;v}_K?p6E|V;~jy^DWA~^m`&;HIkof2E} zm#F|T*R3b6*1jH#Cn%e`LQ=U~HXjfZc|zCERFPi+ShpM^e|^jrF`3JVu5~0FIP=Z6 zECVf@pb5j%kHsV|p*4ydGT_Os=1}F*lYfG4i-~+IN?=%q&W$pYaUYP*t^7AN;NR5q z#Jdy)Xvw@8&T@eEfO~(CN%z6MazHL%?xmZe0!TRnxtYTr50+`)TW_~i1vPBYR_n688Oyz5x-1Q_X~cm2=lsXzrI z)5Dg#7fDj&rb{Afq-)02kbsK7kg;kD8ZlANsNOiE#|OLa;!JSkgP7(}H=9o2xGA|!T$DC*36!%!)wBIpdWYy>EowX z?ze4Ov$TE{_q50DvYH*Rw8X^;;)PTqKzHqm++FGh-sjAkFOEXj)MvO;q6U9^(oQ%X z5b)r6bj2!XvEn)vCuHAtyvv&n(9;SB;>z}!&r2U_gHN()1R*`XB0HR7B&Z|QNy>9rI`lAE%mSA_%wVs&G z;uV=7Y6rUuJ}!iLj4W)doMh;;@2+;f`ex7SUdg zxkOfT;Xg3SI8mIotcA;{EmEofp6dSv!2bWot2ASTl}YDdvs0=oC9Lw5?Wu6?)P0L< z%kthx=DvdJ;5;*_gffB4-V4jxWBr{THqs_*To~sr)Hr=t3?UWdkjB}lcDJ}@asci% zOj#BlkwWRM=Ed%e#`OoImeqFSb(Cu=s&V0VUE-8-7Ek@JZDPfCL8!u^1JhFpJ_-k8S4Q46 z+~}{r>dE1bLQER)$p2JhPxbBNBdo z;y(ueHX|PS*4o|KZ1imKo%U2>$RI)G} z$-k{Sw4Jm0-bfvS&?WtPamI@C9MMq8R&Rd^vyiwgmeMvLaGjxK44B0D)}08Vxegkd znIQ!=*&m;c-m4?pq#EGCbA6y+{MG~p*U?~jWo^H%S5OECn|tb7*|DXkWW+p+xkpqG zTQ-ZMWOX^k5CtX`ARj@0>k&>!TJPk`Jw~&vk6zpat2m@fOg=O2N-{EA>DHAe6SKT*$Vs~F`VMvH9Efqf2eFJ3IA}0+1 z4W9zNZ*I5@YDyL~J$K<6WBFfjqYzp~!Zw#0Ou82f8M!OZY(s0bSFxwV9VK*aR*XMv z8o2K%p7X7>A!%tc`c1i*C}8+H-Afy?HJA1C%GYp|UyeXV*LC^Pw6h38ILG8m;EGeQ zfpbjwVoU+<)xy|}CyXz_;%G(_eJo@cgegtn5PwC5hG1XNVY&Nu1;9t&v)Xm%roYVm#9_u!` z3*{%P2ll-RN%+!}2Nb!N1z!O%Vaczl8-|r(k%X!qHNGoQDHlJGzk%SZsI8pbJnZa1 zrEs`m?zI8$Pg3U}YKqyg4v+hR?#QiVh05B{yR7|*Vn^KO$EtU%LNFF0ro&@B)Jf`A#jRzgcl zNSw%N@)@AgEXwPm=vT2gJ;!>W;ge(}SbynnE5j(`gYs`6d@=j}+P~twW|eL0ua!dP zZ8)sx(~p)XDEmR?wpGdLpe<>TPAla47ctAUHo{PdR21%-C2CGN>(!RO6n`nJ#{&Sy z2hP}_;whhTRR@5Ug-C*cX70f%hdw!T4He8iJ7stAGyIMv&%F8r*s(z#T?3IH6r>0Z?X=ry1Wu*unD2Qzo8@F zub3Y=_>XL@cahx z*7G5QVNZq^zFxBmAPzeHmD;&Wj5&nmjlJNwk%+#vy&(=hr!ss{ zlF&6yVhT#H%h~bR2qby_nE_eH8*x5u?NUX5of-AGDpQpM{J2%zrxAEhYAf6dRGwqib4n`O0Q{<@`D|* zM5DervNR!zNecgoNuZ+iU!Sa$xVhg}VP?*wLu1$LMCM9u_RM=xZPx#FKdZBCx(mu5 z2&w?;7liVjP|c>|yw1>#V&5Gg;0~ zMYg6cz*kpr_+RS>sVaQ`P=Z`ar*i0zbEEZ~ccUKSB^Ur3m<=(??%Vh9Bj zy&JDGl%NGlNW2@+m6F7-JxB#Z=GQJB?;`>0*)R5jYfDHSelOC6fx15~cyG;`WOjBk zKaaJ7n-p8kc}CgYvAY}$#<)F3h(nvD#is3uaM4K*n~8V1+4sC$_*4|sgl-)fdm!C4 z)n#0z*pMq&h5SvdixrS&o{H)wpJ*WPg$6_KIE2@cE-?2#m}~H6#n=Xw@B5WGw9O-Y z-F0SFA47MY>Ac5GHsRu0Xy@3A+yWi$Q~>u0Sa9aLpGnd`i6kYcili7>1d^HUc4__LPCG+>h1bZ2~S~PpYEdETcTZdqcy5`byznlsw<3 zhg}U$#IR<5^#lQr0`vq|Jc>}%^8=T!Tw9y z3>l9O-aVpdc+nki19TtpFn|*5A4>V%+hmqcGE_m9Pcqog2LuM*tMu1A1JTVqZ7n{{ zklTdotcnF|x1HPo1d-pwwW;MY^|G6B2*;cm{&b#}iAC-=6GP}hGEQeuhj3)Jbe#fi z&tHBOggNKEkd=E$2iEh`=xy3WQ9?YzcZ3fQ@87%i@4LgdWI{VkCQWA4KVX&d6}G3d zuE2e|Z36D)ii6x{Glf{0gM91XO;^F5$FE3L9ufcsHK)(Vfpo_wo`Q$Kj z6NQFw==C_ks=d3#6g0^cLU*_+db(FVOV`&uE6WV1YwQy1RSw;|j{E5dRQYH?#Pu^H zz4N7hZ=}Tmq!U`ozkXS;|E3i8-_W<5=>K;$UQ`|Xze*V{SfVU*%K4jhfwxFSb%cL8 z0TOse8~S6s_4HLSTmgr{maDU&_*rfB5cWF~8~=%T37)urJLfXgoSaIcuc6&O)N{@25mbN4(NwuvO~xtnlv}--<|S z#qAidUv5{F0a&&G5+g2rZ#P7|P{e`#V8pd}S8N=8ua$V=0LsIq!WJuyIdsY3Qn z5#Yd|%9JL^E8Fimsp6>wecclnh|uJ?=cKS{qnU| zR)*SP9q)>=63V2GZ%t6VBw;A?{ZT$d&OEsQ+?lrLD!(HBU$#lf*$xYI%yy^Xz_r)@ z-scO*yc@4&TEVWbF7!r`l+%Ufc;bSYLVK*xOo5w2m)5astR}TyL~K5`;%?h~Yc6wL z*ERP|bzvEKx5L(u&yzyUFT<23p=f_MZ3$k9lQXnE3#@A+XBx;))Hx>_H!bW`}g#qdN5I>5$=% zxewR=D&VPGO6b}zKm1R8^WV_(e;xGvcQh}8$$wPT7!3i3unbSThz+Nwm_#EU)KH}Y zHCx4HZSHHgptWlo4spGJqesh)nnE*_^(2uoadip%vv1X}CXew8NR~~ikqdR8Z^(uA z&VIvtamWox0uFoh>JNTwJ=j^sSpTc;f!I$z*$tdL9mxQd3U>ZF1c2C0=AI}B`7t{5 z7&x{2&&3en;0HFtf2Z6xaqf<~h)A{3o_ojS?C(_l`v?K>G_Ji02llPA2F{H{pH}|0 zp&}A6ivL0J2zw=<26j}a`=<}S0YG>$tRha3ah!8JMQQTN;f>^CQapH%H455AT?hIw z8x8c0f=-Z&rEG!YrkmpeG3E{B@UD+*L}zgi-9TX$lX_l4$kD+dP~jNRSggHfczgozf@;^1o$V02yQke9MD3aI8TfW6~s0UPU^nrX|V~)`1Q){scaqp&Rml z{_7aeCcO>_g`TaTBGTl-^$|7MxWPiT2peU|a-4mg4R|6z5KXv|HsfxPb?AnQU%vYW zsB2FmAZXi#&dArwV1C;TsA{=v9IpWP{7F~j+|zQt`9bv+bvw!{JCSW5f$0KPFhW)GMW;OL2B#l5PU$ZLRE8zz5rjj z==l+pv0awj(F!}D*w=hxoMMMAzdPIO%D0iko|o$uv-ZBm9&T2bjC?m8;>&)fzaaSrHKUfx`N*(9h=>?$=Ps{?6dgsM)BBM#fc#p%e?ltdg0- zbJlo+D+~KNyN;2ujJL2o-w1@?LbB3r)$F$qUZ10@*JaNjhaH`*c2UMo z2?r(;C$_*!ZMuFa>g1$y$WKBY%@lls;$abw=2gC$@bYRaq2u?6qBE>;zdr(H{iX-I zX@eQN7Gv7akYKg=YPloZ*LtKIz?ubr@fn z5&f++FBmsa+T7&`9O=RK*Z}-N?oaoCO-r5gvS@u>qC0;Mza?{chiwJNhEbR2g~h99 z1eBk@i+0|lT{v6+fc}2EM}J#4Cb%3*yy75yrX`IFICLNpicgHEduVv^t7fb|9P;kH zoILxIh@oDtK~<;5`I;D&LI$NZbbJvu5vSL~aYYFsVf@*66zlK^9_($`?JE&$^KPfk_}6uDQSw1|>aJU;y`aiU!czGiKfREa zw)W}UFNd`AuhbYFeK%A)hz8)!A-X?;#1DocbaS4K#(rTP{W`6+?}{*_@<=W`>u|`_ zIKN*11n$)BhNpt8IQ;BLkPNjP0KVH(Ec{gS*&yrm*8AaWO1uPVA?U`Eir|7!J1zIV zp?gEg;be8S%f@4OenZO`6klT3uq`Wp(=Z2C9UpPrT*1c-6(4f#27v~%eT$hGxRU-n zW6_k_dwf2n^uB@d`}(8Ow@>KUmAvg*EVfBv4t|PwS_queSW_jP#OcoJbVYQH=7rcc z64lA-UTmiss(08Ff?Jr_)}QiLZKjO#!b3{C)t()>bxz~?#w)&MDCJbn!ifm5)?v~#EVW8~Oee}Ek4xpw)nrTrnN`;! zx~@nG_esdTp!<5?B9&3{o@~mv==!+DmF;mf5udP|9Ges3$Qtz5md4#WO*>cQnP#7AvZA0!%=_X^u(?*vOB$VwEofAmBe#^HD!c-xV33Y%U7 ztX3g~;&X^iZ6GIo7SQwi>vPvnOj>$B?S1Haea!VwJ@nO$qd_GpGUF$MeWv%qr57Y3 zGj!W0W~bXEZa>m-AsVF#t9R{CgSt$Ij7niVNoB&;Pz*em47Hb6N6R$iNVtVRk#BR! z-3Fi}hSX8gELgRcyj5FicD@a7BaKVSy=8T^x%Q!rL4UA&*B_mAE#5H;HK`OA6>0M$ z+AzJ(7-yXdM$yu}%@YBZ3hZO!UJUAYTXNF8P1u2s@E7Cf+34qYzTUPRO8K-(`1YJs z=zP|@8hTi!S6tfG6C=U7nr_y$Zi1Xf0a+c(KHQxF96y)A4~O$o1`eJq>TTF`8BvX+ zw{7(NW_5|Yb*U|`C$C*FF@l*fcioYd>sIUSVsBp1QdTphnt4}_?&(=<@}wJ%c>U2m zz0(rWm8doHYRP9Tb}t^M4^CfAR`0~!;judTPa5zYv97maA%+mRtg$& zu5ybxS#zGVp0UOJ!3MJWCbnAOenW)Q;bWL@xDRfMl$0*Th^N6?!y6ycJ=kQWdCy+7 zxpSFvqF=#FM;p`XV!Ru7 zr>O5v(BQG7@-L;Kp}IwxUP zh^1{=JDsdPrP94L=(gw9OY(im_=;u5KyB?8q6p>+Tog*@-J@7H-b)CnwVFE8DS9h?cyz+!%V*wBQ}Ny>;W#!sL*u4)r$p%^ws8ntAM=;C~3&B!;O$C5uePNzHNeJ>&T zFi$-h-SkM|jvtGq37sTq?88n`AJ!l1?hLMg;3FLW1Xkl#BY2}Sv9C_Hg z@}{OLUu!6Q_40c=`Vbw@7~^$kz?eki8x^qeX>3je><8&4W2inF*=}=(N^peTtl_^*8b`L%xzQg4+5-{tZ&v)uKLA2_!`3ht%QI2#&P2$`*V>u}9F>lO5}6AaN2 z%pP0p4R_Lx8qz{xkL7luwSDx^@U7i}fV-e1__0d0q*d#D)y;!22W**hYS?gZ1JfCI zGWvid4@ZAwQ)v=peBD%P>^lfpmtj#?eUd;Ze&Fe1wF%L_44C%*rR=o4#Vjg#0106=@}Hgt<}Fn@7cm`IEuzAQiQq2t-Md&7BX4<4!mI zUgIHdlpQlCvX0e+8t3HHsC+1=qi`&Wmd>O}%)YT=2Sm4bD5~Y_ehb=~bVKA^(qlF3 z{-b$L7ZuJVP47YNz^@%4norh6AZ`QE*QK4m*CCJ3KZQlT&e9?@R>}8UnuV9T^xi;9 zWO|>#l^(<*kLsii8E4PRl^lXJlX_V6{tTtwY%e%_PUR5W=>1% zcDSQ+U4(na)R4VRV+2Ox5qH=r>-~_E;QL;9gf-PsaFdmvK`sAifHy0S`^HtKy1}wB zL&w$b?H)7&u=uc}&z=TKV1{3zi?Dv&EB>Rsymz=Mf88?3@r9)dE-NtPy+ z$iMP8cu(gl%D@eA$2yCHcG!le#q;SYC1-p4#=TaVrb7^QeCrNCSX zg^&QLe#W;$)UsKLNe8BT??i4~J-nu38_+C8V6O?FSTRY=y z0i@9JfOO3tII9z*r05L&Ge;X2Ra`a^>HW!mi-#Kp#&yjCr&Qvpt`GJ$###sEYY+*V z7fvxetJ%H^(N-_jWD#!EA=8q`~D)|L0xHkPA?gETwrv0 z|M|50#lZQ?EnKI4!WrAW!(>Xo?QOLk#AF3PC^l_XJ0C;O ze&}8YHq5rv3m_J-cIcRJE;|u#@vd(^kAcj}M$81lxV*`*Lwfiiz`0(IM+h-4e>}yX zC)SsT=(3{2*Lo8%(svmc66akh_5J&Z$SP+CUm+`gWZKdluaqhboZR8=lgG2Dcr}Q+ z9`Wqebb1XC4r`%#AN%#UOFvM?myKY7&nYrF4%$vAY>V{Od8EMErHEAAmd43GvUDJZ z6o6xd+3n_aGg2R#BMbzI&vHF zrMaM>qVn$XI!{*T;y&y|Z=9CPAnjcWe_ni^VZCqf^tu8U~);Ap?5(o8yh6ggW9MS^PhV?&d8!;KZYS!AbJ8hNA zQi;4C7y9{Ai9l1nC_SZ>ps`o5YpQ+v6xjIuRpC^*!Z}>IHW>(U?iG2v9hFcZRG;#C z3}zR>t+P>GO=t^=m^hN$aUJ>7U`cj;UPw#qUV<^rbuD?-x2l}LY>V0M8%D_a%l&x) zRQ*z;oxz1{f&zKfe&=4W9&Kq6+vRiXb=tk@<7%rOL^S5$jNa6o#j9O z=~Ds(1bY0facPz+SPugE2 Date: Tue, 17 Sep 2024 22:05:28 -0400 Subject: [PATCH 09/12] Rename ComparsionChart.png to ComparisonChart.png --- img/{ComparsionChart.png => ComparisonChart.png} | Bin 1 file changed, 0 insertions(+), 0 deletions(-) rename img/{ComparsionChart.png => ComparisonChart.png} (100%) diff --git a/img/ComparsionChart.png b/img/ComparisonChart.png similarity index 100% rename from img/ComparsionChart.png rename to img/ComparisonChart.png From 5867c59b8c7bd856ee35bb20acf60bc0ee3a75ae Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 22:08:12 -0400 Subject: [PATCH 10/12] readme update --- README.md | 2 -- src/main.cpp | 2 +- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/README.md b/README.md index 5b207b0..a2f3df1 100644 --- a/README.md +++ b/README.md @@ -59,11 +59,9 @@ This repository compares the implementation of scan or prefix scan algorithm on [ 3 1 1 2 3 2 1 1 3 3 2 3 3 ... 3 1 ] passed ==== work-efficient compact, power-of-two ==== -zeropadded = 256 elapsed time: 0.077728ms (CUDA Measured) passed ==== work-efficient compact, non-power-of-two ==== -zeropadded = 256 elapsed time: 0.083264ms (CUDA Measured) passed ``` diff --git a/src/main.cpp b/src/main.cpp index 896ac2b..600c29f 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 << 20; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; From 84cf726b2223c5a0a87609bb1e0a13dfd3fb1452 Mon Sep 17 00:00:00 2001 From: manvi27 Date: Tue, 17 Sep 2024 23:19:30 -0400 Subject: [PATCH 11/12] code correction --- stream_compaction/naive.cu | 41 +++++++++++++++++++++++--------------- 1 file changed, 25 insertions(+), 16 deletions(-) diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 139f7d8..bac7481 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -14,27 +14,27 @@ namespace StreamCompaction { return timer; } // TODO: __global__ - __global__ void scan_global(int n, int *odata, int *idata, int *temp) + __global__ void scan_global(int n, int *odata, int *idata, int *temp,int offset,int pout) { int thid = threadIdx.x + (blockIdx.x * blockDim.x); - int pout = 0, pin = 1; // Load input into global memory. // This is exclusive scan, so shift right by one // and set first element to 0 - - temp[pout * n + thid] = (thid > 0) ? idata[thid - 1] : 0; + int pin = 1 - pout; + if (thid >= offset) + temp[pout * n + thid] = temp[pin * n + thid - offset] + temp[pin* n + thid]; + else + temp[pout * n + thid] = temp[pin * n + thid]; __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] + temp[pin* n + thid]; - else - temp[pout * n + thid] = temp[pin * n + thid]; - __syncthreads(); - } odata[thid] = temp[pout * n + thid]; // write output } + + __global__ void shiftInput(int* idata, int* shifted_input) + { + int thid = threadIdx.x + (blockIdx.x * blockDim.x); + shifted_input[thid] = (thid > 0) ? idata[thid - 1] : 0; + __syncthreads(); + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ @@ -47,11 +47,20 @@ namespace StreamCompaction { result = cudaMalloc((void**)&g_odata,zeropadded_n*sizeof(int)); result = cudaMalloc((void**)&temp,2 * zeropadded_n * sizeof(int)); cudaMemcpy(g_idata,idata,sizeof(int)*n,cudaMemcpyHostToDevice); - int threadsPerBlock = 256; + int threadsPerBlock = 1024; int blocksPerGrid = (zeropadded_n + threadsPerBlock - 1) / threadsPerBlock; - scan_global<<>>(zeropadded_n,g_odata,g_idata,temp); + + int offset = 1; + int pout = 0; + shiftInput<<>>(g_idata, temp); + for (int i = 0; i < ilog2ceil(n); i++) { + pout = 1 - pout; + scan_global<<>>(zeropadded_n,g_odata,g_idata,temp,offset,pout); + offset *= 2; + } + cudaMemcpy(odata,g_odata,sizeof(int)*n,cudaMemcpyDeviceToHost); - for (int i = 0; i < n; i++) + for (int i = 0; i < 257; i++) { //printf("%d %d\n", idata[i], odata[i]); } From 82995d86211bc93d6235f2569b38c3a81e2eb3f5 Mon Sep 17 00:00:00 2001 From: manvi27 Date: Wed, 18 Sep 2024 02:55:16 -0400 Subject: [PATCH 12/12] final change --- src/main.cpp | 2 +- stream_compaction/efficient.cu | 85 ++++++++++++++++++---------------- stream_compaction/thrust.cu | 7 ++- 3 files changed, 48 insertions(+), 46 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 600c29f..2fce82c 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 20; // 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]; diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index e019f99..fdaa8cf 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -13,47 +13,40 @@ namespace StreamCompaction { return timer; } - __global__ void scan_global(int n, int *odata, int *idata) + __global__ void scan_upstream(int n, int *idata,int offset) { int thid = threadIdx.x + (blockIdx.x * blockDim.x); //load data in the global memory + int ai = offset * (2*thid + 1) - 1; + int bi = offset * (2*thid + 2) - 1; - for(int offset = 1; offset < n; offset*=2) - { - int ai = offset * (2*thid + 1) - 1; - int bi = offset * (2*thid + 2) - 1; - - if (ai < n && bi < n) - { - - idata[bi] += idata[ai]; - } - if(thid == 0) - idata[n-1] =0; - __syncthreads(); - - } - - for (int offset = (n/2); offset >= 1; offset /= 2) // traverse down tree & build scan + if (ai < n && bi < n) { - int ai = offset * (2 * thid + 1) - 1; - int bi = offset * (2 * thid + 2) - 1; - if ((ai < n) && (bi < n)) { - float t = idata[ai]; - idata[ai] = idata[bi]; - idata[bi] += t; - } - - __syncthreads(); - + idata[bi] += idata[ai]; } - if (thid < n/2) - odata[2 * thid] = idata[2 * thid]; // write results to device memory + if(thid == 0) + idata[n-1] =0; __syncthreads(); - if(thid < n/2) - odata[2 * thid + 1] = idata[2 * thid + 1]; + + } + + + + __global__ void scan_downstream(int n, int* idata,int offset) + { + int thid = threadIdx.x + (blockIdx.x * blockDim.x); + //load data in the global memory + int ai = offset * (2 * thid + 1) - 1; + int bi = offset * (2 * thid + 2) - 1; + if ((ai < n) && (bi < n)) { + float t = idata[ai]; + idata[ai] = idata[bi]; + idata[bi] += t; + } + __syncthreads(); + } /** @@ -61,24 +54,32 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startGpuTimer(); - // TODO - int *g_odata; + int *g_idata; int zeropadded_n = pow(2, ilog2ceil(n)); - cudaMalloc((void**)&g_odata, zeropadded_n * sizeof(int)); cudaMalloc((void**)&g_idata,zeropadded_n * sizeof(int)); cudaMemcpy(g_idata,idata,n*sizeof(int),cudaMemcpyHostToDevice); - int threadsPerBlock = 256; + int threadsPerBlock = 1024; int blocksPerGrid = ((zeropadded_n /2) + threadsPerBlock - 1) / threadsPerBlock; - scan_global << > > (zeropadded_n, g_odata, g_idata); - - cudaMemcpy(odata,g_odata,sizeof(int)*n,cudaMemcpyDeviceToHost); + int offset = 1; + for (int i = 0; i < ilog2ceil(n); i++) + { + scan_upstream << > > (zeropadded_n, g_idata,offset); + offset *= 2; + } + cudaDeviceSynchronize(); + offset = zeropadded_n / 2; + for (int i = 0; i < ilog2ceil(n); i++) + { + scan_downstream << > > (zeropadded_n, g_idata,offset); + offset /= 2; + } + cudaMemcpy(odata,g_idata,sizeof(int)*n,cudaMemcpyDeviceToHost); cudaFree(g_idata); - cudaFree(g_odata); timer().endGpuTimer(); } @@ -136,9 +137,11 @@ namespace StreamCompaction { StreamCompaction::Common::kernScatter<<>>(zeropadded_n, g_odata,g_idata, g_bools, indices); cudaMemcpy(odata,g_odata,zeropadded_n*sizeof(int),cudaMemcpyDeviceToHost); - cudaFree(bools); + cudaFree(g_bools); cudaFree(indices); cudaFree(g_odata); + cudaFree(g_idata); + free(bools); timer().endGpuTimer(); return temp_array[zeropadded_n-1]; diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index d907c7f..1b4db06 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -24,12 +24,11 @@ namespace StreamCompaction { thrust::host_vector thrust_idata(idata, idata + n); thrust::host_vector thrust_odata(odata, odata + n); - thrust::device_vector thrust_dev_idata = thrust_idata; - thrust::device_vector thrust_dev_odata = thrust_odata; + thrust::device_vector thrust_dev_idata(idata,idata+n); + thrust::device_vector thrust_dev_odata(odata,odata+n); thrust::exclusive_scan(thrust_dev_idata.begin(), thrust_dev_idata.end(), thrust_dev_odata.begin()); - thrust::copy(thrust_dev_odata.begin(), thrust_dev_odata.end(), thrust_odata.begin()); - std::memcpy(odata, thrust_odata.data(), n * sizeof(int)); + thrust::copy(thrust_dev_odata.begin(), thrust_dev_odata.end(), odata); timer().endGpuTimer(); } }