From 5ff0183b067c417ce2ea22847ee3baa0f6e63a54 Mon Sep 17 00:00:00 2001 From: Zijing Date: Thu, 17 Sep 2020 16:56:58 -0400 Subject: [PATCH 1/5] Add native scan --- stream_compaction/cpu.cu | 36 ++++++++++++++++----- stream_compaction/naive.cu | 64 ++++++++++++++++++++++++++++++++++++-- 2 files changed, 91 insertions(+), 9 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..565e19f 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -18,9 +18,12 @@ namespace StreamCompaction { * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); - // TODO - timer().endCpuTimer(); + //timer().startCpuTimer(); + odata[0] = 0; + for (int i = 0; i < n - 1; i++) { + odata[i + 1] = odata[i] + idata[i]; + } + //timer().endCpuTimer(); } /** @@ -30,9 +33,14 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int o = 0; + for (int i = 0; i < n; i++) { + if (idata[i]) { + odata[o++] = idata[i]; + } + } timer().endCpuTimer(); - return -1; + return o - 1; } /** @@ -42,9 +50,23 @@ namespace StreamCompaction { */ int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int* map = new int[n]; + for (int i = 0; i < n; i++) { + map[i] = idata[i] ? 1 : 0; + } + int* sout = new int[n]; + scan(n, sout, map); + int o = 0; + for (int i = 0; i < n; i++) { + if (map[i] != 0) { + odata[sout[i]] = idata[i]; + o++; + } + } timer().endCpuTimer(); - return -1; + delete[] map; + delete[] sout; + return o - 1; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..d9569f6 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -11,14 +11,74 @@ namespace StreamCompaction { static PerformanceTimer timer; return timer; } - // TODO: __global__ + + int* dev_bufferA; + int* dev_bufferB; + int numObjects; + + __global__ void kernNaiveScan(int N, int* A, int* B, int d) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= N) { + return; + } + int temp = powf(2, d - 1); + if (index < temp) { + if (d % 2 == 0) { + B[index] = A[index]; + } + else { + A[index] = B[index]; + } + return; + } + if (d % 2 == 0) { + B[index] = A[index - temp] + A[index]; + } + else { + A[index] = B[index - temp] + B[index]; + } + } + + void initSimulation(int N, const int* B) { + numObjects = N; + cudaMalloc((void**)&dev_bufferA, N * sizeof(int)); + cudaMalloc((void**)&dev_bufferB, N * sizeof(int)); + int a = B[0]; + cudaMemcpy(dev_bufferA, &a, sizeof(int), cudaMemcpyHostToDevice); + cudaMemcpy(dev_bufferB, B, N * sizeof(int), cudaMemcpyHostToDevice); + + cudaDeviceSynchronize(); + } + + void endSimulation() { + cudaFree(dev_bufferA); + cudaFree(dev_bufferB); + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { + + initSimulation(n, idata); + const int blockSize = 64; timer().startGpuTimer(); - // TODO + dim3 numBoidBlocks((n + blockSize - 1) / blockSize); + int dmax = ilog2ceil(n); + + for (int i = 1; i <= dmax; i++) { + kernNaiveScan << > > (n, dev_bufferA, dev_bufferB, i); + cudaDeviceSynchronize(); + + } + if (dmax % 2 == 0) { + cudaMemcpy(odata + 1, dev_bufferB, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); + } + else { + cudaMemcpy(odata + 1, dev_bufferA, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); + } + odata[0] = 0; + endSimulation(); timer().endGpuTimer(); } } From 4953cc88120418d39b5ecd43872600508fbf4aa7 Mon Sep 17 00:00:00 2001 From: Zijing Date: Mon, 21 Sep 2020 17:10:10 -0400 Subject: [PATCH 2/5] Add efficient scan, compact, thrust scan --- stream_compaction/common.cu | 17 +++++-- stream_compaction/cpu.cu | 24 ++++++---- stream_compaction/efficient.cu | 84 ++++++++++++++++++++++++++++++++-- stream_compaction/naive.cu | 10 +--- stream_compaction/thrust.cu | 22 +++++++-- 5 files changed, 130 insertions(+), 27 deletions(-) diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..f15f11a 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -23,7 +23,13 @@ namespace StreamCompaction { * which map to 0 will be removed, and elements which map to 1 will be kept. */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { - // TODO + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + if (idata[index] != 0) { + bools[index] = 1; + } } /** @@ -32,8 +38,13 @@ namespace StreamCompaction { */ __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { - // TODO + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= n) { + return; + } + if (bools[index] == 1) { + odata[indices[index]] = idata[index]; + } } - } } diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 565e19f..56e3104 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -12,18 +12,22 @@ namespace StreamCompaction { return timer; } + void cpu_scan(int n, int* odata, const int* idata) { + odata[0] = 0; + for (int i = 0; i < n - 1; i++) { + odata[i + 1] = odata[i] + idata[i]; + } + } + /** * CPU scan (prefix sum). * For performance analysis, this is supposed to be a simple for loop. * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - //timer().startCpuTimer(); - odata[0] = 0; - for (int i = 0; i < n - 1; i++) { - odata[i + 1] = odata[i] + idata[i]; - } - //timer().endCpuTimer(); + timer().startCpuTimer(); + cpu_scan(n, odata, idata); + timer().endCpuTimer(); } /** @@ -40,7 +44,7 @@ namespace StreamCompaction { } } timer().endCpuTimer(); - return o - 1; + return o; } /** @@ -55,7 +59,7 @@ namespace StreamCompaction { map[i] = idata[i] ? 1 : 0; } int* sout = new int[n]; - scan(n, sout, map); + cpu_scan(n, sout, map); int o = 0; for (int i = 0; i < n; i++) { if (map[i] != 0) { @@ -63,10 +67,10 @@ namespace StreamCompaction { o++; } } - timer().endCpuTimer(); delete[] map; delete[] sout; - return o - 1; + timer().endCpuTimer(); + return o; } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..906cf72 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,12 +12,64 @@ namespace StreamCompaction { return timer; } + int numObjects; + int* dev_buffer; + int* dev_booleanBuffer; + int* dev_scanBuffer; + int* dev_idata; + int* dev_odata; + + __global__ void kernUpSweep(int N, int offset, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= N) { + return; + } + //int offset = powf(2, d); + if (index % (offset * 2) == 0) { + data[index + offset * 2 - 1] += data[index + offset - 1]; + } + } + + __global__ void kernDownSweep(int N, int offset, int* data) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + if (index >= N) { + return; + } + //int offset = powf(2, d); + if (index % (offset * 2) == 0) { + int t = data[index + offset - 1]; + data[index + offset - 1] = data[index + offset * 2 - 1]; + data[index + offset * 2 - 1] += t; + } + } + + void efficient_scan(int n, int* odata, const int* idata) { + int dmax = ilog2ceil(n); + numObjects = powf(2, dmax); + cudaMalloc((void**)&dev_buffer, numObjects * sizeof(int)); + cudaMemcpy(dev_buffer, idata, n * sizeof(int), cudaMemcpyHostToDevice); + + const int blockSize = 256; + dim3 numBlocks((numObjects + blockSize - 1) / blockSize); + + for (int i = 0; i < dmax; i++) { + kernUpSweep << > > (numObjects, int(powf(2, i)), dev_buffer); + } + cudaMemset(dev_buffer + numObjects - 1, 0, sizeof(int)); + for (int i = dmax - 1; i >= 0; i--) { + kernDownSweep << > > (numObjects, int(powf(2, i)), dev_buffer); + } + cudaMemcpy(odata, dev_buffer, n * sizeof(int), cudaMemcpyDeviceToHost); + + cudaFree(dev_buffer); + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { timer().startGpuTimer(); - // TODO + efficient_scan(n, odata, idata); timer().endGpuTimer(); } @@ -32,9 +84,35 @@ namespace StreamCompaction { */ int compact(int n, int *odata, const int *idata) { timer().startGpuTimer(); - // TODO + + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + cudaMalloc((void**)&dev_booleanBuffer, n * sizeof(int)); + cudaMalloc((void**)&dev_scanBuffer, n * sizeof(int)); + + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + + const int blockSize = 64; + numObjects = n; + dim3 numBoidBlocks((numObjects + blockSize - 1) / blockSize); + + Common::kernMapToBoolean << > > (n, dev_booleanBuffer, dev_idata); + int* host_boolean = new int[n]; + cudaMemcpy(host_boolean, dev_booleanBuffer, n * sizeof(int), cudaMemcpyDeviceToHost); + efficient_scan(n, odata, host_boolean); + int size = odata[n - 1]; + cudaMemcpy(dev_scanBuffer, odata, n * sizeof(int), cudaMemcpyHostToDevice); + Common::kernScatter << > > (n, dev_odata, dev_idata, dev_booleanBuffer, dev_scanBuffer); + + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + + cudaFree(dev_idata); + cudaFree(dev_odata); + cudaFree(dev_booleanBuffer); + cudaFree(dev_scanBuffer); timer().endGpuTimer(); - return -1; + + return idata[n - 1] ? size + 1 : size; } } } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index d9569f6..f79c936 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -43,11 +43,7 @@ namespace StreamCompaction { numObjects = N; cudaMalloc((void**)&dev_bufferA, N * sizeof(int)); cudaMalloc((void**)&dev_bufferB, N * sizeof(int)); - int a = B[0]; - cudaMemcpy(dev_bufferA, &a, sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_bufferB, B, N * sizeof(int), cudaMemcpyHostToDevice); - - cudaDeviceSynchronize(); } void endSimulation() { @@ -59,16 +55,14 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - - initSimulation(n, idata); - const int blockSize = 64; timer().startGpuTimer(); + initSimulation(n, idata); + const int blockSize = 256; dim3 numBoidBlocks((n + blockSize - 1) / blockSize); int dmax = ilog2ceil(n); for (int i = 1; i <= dmax; i++) { kernNaiveScan << > > (n, dev_bufferA, dev_bufferB, i); - cudaDeviceSynchronize(); } if (dmax % 2 == 0) { diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..9cc3d3d 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -14,14 +14,30 @@ namespace StreamCompaction { static PerformanceTimer timer; return timer; } + + int* dev_idata; + int* dev_odata; + /** * 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()); + + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + + thrust::device_ptr dv_idata_ptr(dev_idata); + thrust::device_ptr dv_odata_ptr(dev_odata); + + const int blockSize = 64; + dim3 numBoidBlocks((n + blockSize - 1) / blockSize); + + thrust::exclusive_scan(dv_idata_ptr, dv_idata_ptr + n, dv_odata_ptr); + + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + timer().endGpuTimer(); } } From 0c15df6a11cf7f9ceb186526ff29b533b8a689dc Mon Sep 17 00:00:00 2001 From: Zijing Date: Mon, 21 Sep 2020 20:16:55 -0400 Subject: [PATCH 3/5] Reorganize --- stream_compaction/efficient.cu | 63 ++++++++++++++++++++-------------- stream_compaction/naive.cu | 7 ++-- stream_compaction/thrust.cu | 5 ++- 3 files changed, 45 insertions(+), 30 deletions(-) diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 906cf72..4212d65 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,7 +12,6 @@ namespace StreamCompaction { return timer; } - int numObjects; int* dev_buffer; int* dev_booleanBuffer; int* dev_scanBuffer; @@ -43,14 +42,20 @@ namespace StreamCompaction { } } - void efficient_scan(int n, int* odata, const int* idata) { + + /** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ + void scan(int n, int *odata, const int *idata) { int dmax = ilog2ceil(n); - numObjects = powf(2, dmax); + int numObjects = powf(2, dmax); cudaMalloc((void**)&dev_buffer, numObjects * sizeof(int)); cudaMemcpy(dev_buffer, idata, n * sizeof(int), cudaMemcpyHostToDevice); const int blockSize = 256; dim3 numBlocks((numObjects + blockSize - 1) / blockSize); + + timer().startGpuTimer(); for (int i = 0; i < dmax; i++) { kernUpSweep << > > (numObjects, int(powf(2, i)), dev_buffer); @@ -59,20 +64,14 @@ namespace StreamCompaction { for (int i = dmax - 1; i >= 0; i--) { kernDownSweep << > > (numObjects, int(powf(2, i)), dev_buffer); } + + timer().endGpuTimer(); + cudaMemcpy(odata, dev_buffer, n * sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_buffer); } - /** - * Performs prefix-sum (aka scan) on idata, storing the result into odata. - */ - void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - efficient_scan(n, odata, idata); - timer().endGpuTimer(); - } - /** * Performs stream compaction on idata, storing the result into odata. * All zeroes are discarded. @@ -83,26 +82,41 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - timer().startGpuTimer(); + int dmax = ilog2ceil(n); + int numObjects = powf(2, dmax); cudaMalloc((void**)&dev_idata, n * sizeof(int)); cudaMalloc((void**)&dev_odata, n * sizeof(int)); cudaMalloc((void**)&dev_booleanBuffer, n * sizeof(int)); - cudaMalloc((void**)&dev_scanBuffer, n * sizeof(int)); + cudaMalloc((void**)&dev_scanBuffer, numObjects * sizeof(int)); cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); const int blockSize = 64; - numObjects = n; - dim3 numBoidBlocks((numObjects + blockSize - 1) / blockSize); - - Common::kernMapToBoolean << > > (n, dev_booleanBuffer, dev_idata); - int* host_boolean = new int[n]; - cudaMemcpy(host_boolean, dev_booleanBuffer, n * sizeof(int), cudaMemcpyDeviceToHost); - efficient_scan(n, odata, host_boolean); - int size = odata[n - 1]; - cudaMemcpy(dev_scanBuffer, odata, n * sizeof(int), cudaMemcpyHostToDevice); - Common::kernScatter << > > (n, dev_odata, dev_idata, dev_booleanBuffer, dev_scanBuffer); + + dim3 numBlocks((n + blockSize - 1) / blockSize); + + timer().startGpuTimer(); + + Common::kernMapToBoolean << > > (n, dev_booleanBuffer, dev_idata); + + cudaMemcpy(dev_scanBuffer, dev_booleanBuffer, n * sizeof(int), cudaMemcpyDeviceToDevice); + + for (int i = 0; i < dmax; i++) { + kernUpSweep << > > (numObjects, int(powf(2, i)), dev_scanBuffer); + } + + cudaMemset(dev_scanBuffer + numObjects - 1, 0, sizeof(int)); + for (int i = dmax - 1; i >= 0; i--) { + kernDownSweep << > > (numObjects, int(powf(2, i)), dev_scanBuffer); + } + + int size = 0; + cudaMemcpy(&size, dev_scanBuffer + n - 1, sizeof(int), cudaMemcpyDeviceToHost); + + Common::kernScatter << > > (n, dev_odata, dev_idata, dev_booleanBuffer, dev_scanBuffer); + + timer().endGpuTimer(); cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); @@ -110,7 +124,6 @@ namespace StreamCompaction { cudaFree(dev_odata); cudaFree(dev_booleanBuffer); cudaFree(dev_scanBuffer); - timer().endGpuTimer(); return idata[n - 1] ? size + 1 : size; } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index f79c936..e14af3b 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -55,16 +55,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(); initSimulation(n, idata); const int blockSize = 256; dim3 numBoidBlocks((n + blockSize - 1) / blockSize); int dmax = ilog2ceil(n); + + timer().startGpuTimer(); for (int i = 1; i <= dmax; i++) { kernNaiveScan << > > (n, dev_bufferA, dev_bufferB, i); } + + timer().endGpuTimer(); + if (dmax % 2 == 0) { cudaMemcpy(odata + 1, dev_bufferB, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); } @@ -73,7 +77,6 @@ namespace StreamCompaction { } odata[0] = 0; endSimulation(); - timer().endGpuTimer(); } } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 9cc3d3d..7de513e 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -22,8 +22,6 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - cudaMalloc((void**)&dev_idata, n * sizeof(int)); cudaMalloc((void**)&dev_odata, n * sizeof(int)); cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); @@ -34,11 +32,12 @@ namespace StreamCompaction { const int blockSize = 64; dim3 numBoidBlocks((n + blockSize - 1) / blockSize); + timer().startGpuTimer(); thrust::exclusive_scan(dv_idata_ptr, dv_idata_ptr + n, dv_odata_ptr); + timer().endGpuTimer(); cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); - timer().endGpuTimer(); } } } From 759199370782fc67ca39b74ec6fbf55859b75ff2 Mon Sep 17 00:00:00 2001 From: Zijing Date: Mon, 21 Sep 2020 23:57:18 -0400 Subject: [PATCH 4/5] Optimize --- stream_compaction/efficient.cu | 9 ++++----- stream_compaction/naive.cu | 29 +++++++---------------------- stream_compaction/thrust.cu | 21 +++++---------------- 3 files changed, 16 insertions(+), 43 deletions(-) diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 4212d65..0c469ba 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -23,7 +23,6 @@ namespace StreamCompaction { if (index >= N) { return; } - //int offset = powf(2, d); if (index % (offset * 2) == 0) { data[index + offset * 2 - 1] += data[index + offset - 1]; } @@ -34,7 +33,6 @@ namespace StreamCompaction { if (index >= N) { return; } - //int offset = powf(2, d); if (index % (offset * 2) == 0) { int t = data[index + offset - 1]; data[index + offset - 1] = data[index + offset * 2 - 1]; @@ -52,14 +50,15 @@ namespace StreamCompaction { cudaMalloc((void**)&dev_buffer, numObjects * sizeof(int)); cudaMemcpy(dev_buffer, idata, n * sizeof(int), cudaMemcpyHostToDevice); - const int blockSize = 256; + const int blockSize = 512; dim3 numBlocks((numObjects + blockSize - 1) / blockSize); - + timer().startGpuTimer(); for (int i = 0; i < dmax; i++) { kernUpSweep << > > (numObjects, int(powf(2, i)), dev_buffer); } + cudaMemset(dev_buffer + numObjects - 1, 0, sizeof(int)); for (int i = dmax - 1; i >= 0; i--) { kernDownSweep << > > (numObjects, int(powf(2, i)), dev_buffer); @@ -92,7 +91,7 @@ namespace StreamCompaction { cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); - const int blockSize = 64; + const int blockSize = 512; dim3 numBlocks((n + blockSize - 1) / blockSize); diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index e14af3b..6256c52 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -16,27 +16,16 @@ namespace StreamCompaction { int* dev_bufferB; int numObjects; - __global__ void kernNaiveScan(int N, int* A, int* B, int d) { + __global__ void kernNaiveScan(int N, int* A, int* B, int temp) { int index = (blockIdx.x * blockDim.x) + threadIdx.x; if (index >= N) { return; } - int temp = powf(2, d - 1); if (index < temp) { - if (d % 2 == 0) { - B[index] = A[index]; - } - else { - A[index] = B[index]; - } + A[index] = B[index]; return; } - if (d % 2 == 0) { - B[index] = A[index - temp] + A[index]; - } - else { - A[index] = B[index - temp] + B[index]; - } + A[index] = B[index - temp] + B[index]; } void initSimulation(int N, const int* B) { @@ -63,18 +52,14 @@ namespace StreamCompaction { timer().startGpuTimer(); for (int i = 1; i <= dmax; i++) { - kernNaiveScan << > > (n, dev_bufferA, dev_bufferB, i); - + kernNaiveScan << > > (n, dev_bufferA, dev_bufferB, int(powf(2, i - 1))); + std::swap(dev_bufferA, dev_bufferB); } timer().endGpuTimer(); - if (dmax % 2 == 0) { - cudaMemcpy(odata + 1, dev_bufferB, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); - } - else { - cudaMemcpy(odata + 1, dev_bufferA, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); - } + cudaMemcpy(odata + 1, dev_bufferB, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); + odata[0] = 0; endSimulation(); } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 7de513e..6226afa 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -14,30 +14,19 @@ namespace StreamCompaction { static PerformanceTimer timer; return timer; } - - int* dev_idata; - int* dev_odata; - /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - cudaMalloc((void**)&dev_idata, n * sizeof(int)); - cudaMalloc((void**)&dev_odata, n * sizeof(int)); - cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); - - thrust::device_ptr dv_idata_ptr(dev_idata); - thrust::device_ptr dv_odata_ptr(dev_odata); - - const int blockSize = 64; - dim3 numBoidBlocks((n + blockSize - 1) / blockSize); + thrust::host_vector host_idata(idata, idata + n); + thrust::device_vector dev_idata(host_idata); + thrust::device_vector dev_odata(n); timer().startGpuTimer(); - thrust::exclusive_scan(dv_idata_ptr, dv_idata_ptr + n, dv_odata_ptr); + thrust::exclusive_scan(dev_idata.begin(), dev_idata.end(), dev_odata.begin()); timer().endGpuTimer(); - cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); - + thrust::copy(dev_odata.begin(), dev_odata.end(), odata); } } } From 4e7ac034c1798c89e2782bba16ad60ba2c295698 Mon Sep 17 00:00:00 2001 From: Zijing Date: Tue, 22 Sep 2020 16:36:42 -0400 Subject: [PATCH 5/5] Update README --- README.md | 108 +++++++++++++++++++++++++++++++++++++++++++++--- img/compact.png | Bin 0 -> 21388 bytes img/scan.png | Bin 0 -> 21828 bytes 3 files changed, 102 insertions(+), 6 deletions(-) create mode 100644 img/compact.png create mode 100644 img/scan.png diff --git a/README.md b/README.md index 0e38ddb..85357bb 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) +* Zijing Peng + - [LinkedIn](https://www.linkedin.com/in/zijing-peng/) + - [personal website](https://zijingpeng.github.io/) +* Tested on: Windows 22, i7-8750H@ 2.22GHz 16GB, NVIDIA GeForce GTX 1060 + +### Summary + +In this project, I've implemented GPU stream compaction in CUDA, from scratch. The goal of stream compaction is to remove `0`s from an array of `int`s. The stream compaction includes three steps: map, scan and scatter. I've also implemented several exclusive scan algorithms, which produces a prefix sum array of a given array. + +All the implementations including: + +- Scan algorithms + - CPU Scan & Stream Compaction + - Naive GPU Scan + - Work-Efficient GPU Scan + - Thrust Scan +- Stream Compaction algorithms + - CPU Stream Compaction + - Work-Efficient GPU Stream Compaction + +### Scan Performance Analysis + +![](/img/scan.png) + + + +The four scan algorithm implementations experience huge performance loss as the size of data increase, especially when the size is over 1M. The CPU scan has good performance when the dataset is small. It runs faster than 3 other scan when the size is smaller than 4K, but after that its performance lose rapidly. When the size increases to 16 M,it is much worse than all other GPU implementations. That is because GPU is designed for thousands of computation in parallel. CPU has limited threads, and there are some optimizations in the OS, so it could run pretty fast when the dataset is small. But when the dataset is super large, it will experience huge performance loss. + +Compared with naive scan, work-efficient scan is not so efficient, it even a little bit worse. As Part 5 discussed, more optimizations could be done to improve the performance of work-efficient scan. + +Among the four implementations, thrust scan is undoubtedly the best. When the dataset is small, the advantage of thrust scan is not so obvious compared to others. However, it the only one that still has good performance when the size of dataset is 16M. I take a look at the Nsight timeline for its execution. I find there are several `cudaDeviceSynchronize` function calls, which means they use shared memory. Moreover, I find that in thrust implementation the kernel is only called once (while the up/down sweep of work-efficient is called 24 times with the same data size). The thrust scan use 40 registers per thread while my implementation only use 16 registers per thread. Thus, I guess it take advantage of shared memory and registers. + +### Stream Compaction Performance Analysis + +![](/img/compact.png) + +The two implementations both experience huge performance loss as the size of data increase, especially when the size. As we have discussed above, the CPU implementation has good performance when the dataset is small and experience huge performance when dataset greatly increase. While the GPU compaction shows great performance when the size is over 1M. + +### Output + +An out put when `SIZE = 256` and `blockSize = 512` . + +``` +**************** +** SCAN TESTS ** +**************** + [ 39 42 3 23 38 47 7 10 32 49 44 21 25 ... 9 0 ] +==== cpu scan, power-of-two ==== + elapsed time: 0.0007ms (std::chrono Measured) + [ 0 39 81 84 107 145 192 199 209 241 290 334 355 ... 6130 6139 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 0.0005ms (std::chrono Measured) + [ 0 39 81 84 107 145 192 199 209 241 290 334 355 ... 6074 6110 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 0.018816ms (CUDA Measured) + passed +==== naive scan, non-power-of-two ==== + elapsed time: 0.018272ms (CUDA Measured) + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 0.081056ms (CUDA Measured) + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 0.044896ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 0.054528ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 0.054112ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 1 2 3 1 0 3 3 0 2 3 2 3 3 ... 3 0 ] +==== cpu compact without scan, power-of-two ==== + [ 1 2 3 1 3 3 2 3 2 3 3 2 2 ... 1 3 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 0.0006ms (std::chrono Measured) + [ 1 2 3 1 3 3 2 3 2 3 3 2 2 ... 3 1 ] + passed +==== cpu compact with scan ==== + elapsed time: 0.004ms (std::chrono Measured) + [ 1 2 3 1 3 3 2 3 2 3 3 2 2 ... 1 3 ] + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 0.092768ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 0.125632ms (CUDA Measured) + passed + +``` + + + + + + -### (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.) diff --git a/img/compact.png b/img/compact.png new file mode 100644 index 0000000000000000000000000000000000000000..b788d090afd24e6e5e3b6bfb911504bae671314a GIT binary patch literal 21388 zcmd3OcQBpZ_b(#iM4}u)B+5A)q7y-aL^(PUEqajXB}5mU=ph`UlZf7XOZ0?9@1mFJ zy%U|gAMgA9{qDUpcjo?cXKrTRnb%|Qwbx#It!ysfxH1+N z4mtQQ8XpIIiI0+cfrW*@k`Whu{uX<^mcU+D-FcDy$b+C_?Zx{mw6AUnet&`Uobc6L z&#U}_IONScI3Yy>xI|*wMc)jcQZnkzDC~1vWUSxF$@H1vw%1ZujO$?*-KPuTDbir* zds_56!pd zMV=VhbVea!FDzUHmiI4Mhgr>aUIZ34z6chM3ND5?^a}KU#Q4%!Fc^NQ`Q*qeuWNehL0rU|I7Y8na#(9@|2l|MQKvG?~oJEYmg2Rcae*Z%pHhue_ zCB1Ml7_UEBc$X_m<2PTD!pB8<`QqYA@m?-2f_{XDPz)3C^z>9$kEvTFlXcZoPzVbd zkP0fZOuo=3iB->=Hg+zP3Z+`!*r=diT3szQGa(fd6LbB$R=i+2{?%r8q0_OnC1`!D zGAuOI1%~UGHD_sPXxJXjewXj{7N7JT92w%C6ZGL%g+IMgv=e1!{YhsdNt7z}+aKOM zKK$duXz9q@T(0DHzMB-&$=!V?qd$@QU8g;Z&EeK7&c(ymA5X*^D+>(L zW}kAJa_4q^`}UF7_U!B|wz<+(oN(foNLF|Vj5nn`kJn=GtzhXYl1ScDOH0c>esEb^ zThMWhi$KrP(sHcAQiYv8(&A_1^PyH90mFpRDu>mkCh>U=_Ckw&nIj(`ALByJH%fIQ zdF37{p@t!q1S*prcj?zN^2)__r?NH2Ki-rIqbVnMNnp{I+2tTt<+!0GWZqVmv6FF? zvp_340Z&z#^|^nkO#rpQCtu2HvMYrn`T6;OZC`Cp)^c~a{2kPO!#hN7KJvCHr1rI` z?(DoN?^oM-*#{y~!uCrTel5rKQKL|7yN2Oc(I0*+QBo$ojIIiUy@B8*2*)NmXvenKE*b= zlQ||omR(u%chrq{W5URCw<2}t7gvAYL}VwUPV^>q9yJLp23CZb$0QEjmuMHQihho* zGqar9-XY-E$)O`RDuL2jTJF>EOfrtyDIaITpz|w^vnB1n=e{W~ zOdjDID%2~?eQl6%+(%YM9b+uVXW6dKc&cH=F)gq2p_t8BcO;|rzOVpSuG_m*$u-<( znP$9nM59&agI&X~GvwFCDiuW5-X)LrX6%+*_4kF<4_N8+sP> zytlf&EY+Z+g8oN^{>T^cG5hDydi=xOf+O}Hu;4TCQs$$n!u)iT@f$fa7j~ZyYY~7qv+p?(H_y z1RHx7?p(br|6FZIw>*O{pK`Q}rp!S-U(2YUdrZU1wj=PhgZ5L{=1cjy;z2c@X8+=x z66^HxD4%^tsZmj!CN!13?30i74%uO7gqGz{RJJuDF)+LD7S9|V&C){?c?NXyqVLMpU~l=qe&E;sGTIW)v`6iJ z{i7n*IPdu;zAW*U?=A^+pB`hAmqeJ**n}BT{8ZI`NI$~&sh=BIMY(iSB}E?kFQn*&Peqqn(@*rcRcgbY*`Co!oGC7 zsxSmnwHO1t{GpUo1Z`|04!b0gC5r{i{@>|+6%lc=NJtp&=aqCy5qG%aS#>4Z311wI z1q;HZ|1jt4WtABwo7L7N**R(wJ%>x)=p~{;O7v)&$*Z9qx5Ms{?j++?Z2SQRFW-Tm zjnD4EObZnMzGh|TSRPSx*@-IH+ea|&Dl$K5QRv3yQF!wld-qWL#MI`=k_^%H{?d!r zp})*?W2G`t=u13olp!7=-}t>RrKjJ0u5|`!Xr(+T{G6M0?fc5B#t35fRdF_?lGdVWIxpR) z!Ve8h`8zJg)y_4=p9Qs=>NV#0q(Sb^1losj77jZp->;`bg zWBmu~uf;6e9w#W^;EXUdu%W8!ShMEisfJbQNvrt~dvxf0wdEDv1OX{Ii_S*!0fSR z`uqFu>3#uGVeO-qqMqQzj%u$x51QIL(s4EKo6A#iU7^tVIRhrkcS$1IWg@|yI{Sul z7f|HdH{tzbIh}G-Wx17=LxHlY{jI*2@w`0@RfS29lZr!0EAd&5&PjYp{QK8&KK3BN z#VOqXGStq*N3kHI=xJ$bhpM903if*6M>Btkd6wWi8R9^?e8p z*ttv`Ft@!Il(Qo>O#U8JuPDtFyB=6E+`4)n@2z^-)6jz$EyjR1msn*TAd&)~=_)G5 zz05*>kOJbYr&y-gmp}H=dHJSEeD**B85(}Zf=4)0aAQI9%t%DgeRK--&|2IISg_U_ z0?JF_m4V#FEId>1Qn~JV`5I38-MGwpv_Q&I*3!=jkkq>ape6Sv8V|~t5bQwFwC0v6 zpwwsqU}~Hb@(PsYV92piDE6H_0%*`KnB7s0!G^>t9Sj&DcRu!m204S-3py1V27QgF!8i!R%cQLoA5H|8*M;H63AImy}w#z3)jTT@-qr^(QGLRvRDv2s_&y zn(V83ES5Swu}o=8o&jh%(<9E_ZYg-Y-S}bjpok%OVe(?t;*EACH|wk+l9l!Ik7t0v z09>AnqT}T0Py9|t!6%;a$u@Mz|C;}0Keq%)?OE_RGtWyUK{`-nqv1Bb$QKRVTNg}s zx5`=D?b@zCDi88~G*1`;jPJd~%U$7<`h{-KqsKLd)8k>|4r2bqHLqR#Lw-Gn1bc+| z(sHCd#KtL4w2+g&D|F_~yZN?1=f=$swEM;_nz7fh@#ioNO~DiinF0HM>T8zsEz}Jy zd^^x3^PG7=Y_<}Ab_|y@o4aF?`R@BA9UpGOS`*v3ZUeCqVJ|oP$$h-jy^>48uANwG zghZi)o9XBQTijBAwuBS8`KG?Qm@S0Fe3zCcul4T)r}hj#ghtR((stLPplt?&!LYg9 zOaN{*k@#@BTbGC4%!oXKG=}>I#CZ^Liwm0$Ce{ySo7Mw^bm2OVW>~s~h_tkydh6w} z<9p{d_g#g-WagxkNufDpxFs}lLtYrn^}-0q^&&{1DFkqML`yjO9(2CO@D@=vohxXx zu}c*DzsrCRi5QGHA-MnqXIvje2=$lm5W~6qb6<->O5_WDnDO;R*%tl7NF1&$@#Qjc zx<&+$GM?h@kjPQoL!r)n{AeKk>Of(_G^&32`gYi}X+!JI+=D(mLL#?m+4dl1e#z+V z*3`$#h%(GD7=etNSQz^769x;qc$s_+VvPppTTr^iQ|Z9wViaUnO$vyI#rDZ)^o+3X z5@dB=0U}m_RZ{qj4Vw%!W1F-^o=qnUgE3zW^Qh15Yg677&_|<0{^0;$v-mlk9@T;(9_d1F)?}brmzxa z1~I6450j};=4J2#{ERO{Q%y(tLSmG7|@5G;d z^P%DlX?p!%t94~TQe}5<-n>R1cUEMW&|0(Af+HtXzXsZl5eDP{wp{!|Q@c8w|xdSob#bu92A zfCPuZvMn(@A~Q9z7@SHdyv2x3Coh~>is(@X&|6zEL}9T#KR&FrQ>je^2vQ0*#~aNt zE*+TKOl1u0=Nl3;q)|oX>pQo=WP|`p3wC=)j=~Dn)N4Rk3WPRJ++x9|Lx@J(+C_2` z112SK_(RhP9>Qm4?gz*NCav)29z}D|B9Y(t6)Ew-KIy!vn*03j!fdu<8ySK89^%?k z$;E_WFymBNH-LAt4i^bESpM1Ou9O&D*WHZ=N27qP@nF3Jm833bUzE6-bZdpm7|=1!C|KsWVvKA%RDPU5Qcx zV(21iac9$M3MXD96ube%up^2%D7OCwgS9pnw=n@REaJ}zMRUwZ2igTt;{h=oAuuNz zRpUSmvl>RXfEes>ZE58wY*0<#8VP-X7^sMu&}=$ZKn&J-+_*pt{5TSU=>!7snep%d zQ6L6F{9DLqj)zDj|06X@93TcNZ@Q+wtJE-?jh03P5W_Xx_r!8B2^b6&RqhRd57xmV zqWa74LVdqF%)Pn^#Bc@YYEXK+inKH#p}Ie?{jU;V5XJVQQs{AkogzFS25j$lO?~ss zDAct8Rxu!kaXLkeM%5R3q!qQZ#ueb$KZ)4rFYnxfN9;F}odEc##J_bfnqxycQ2QF~ z;H7;lCTdhY_VYD7v;OG;Ho=bbGblY+5I!^bJ%AT%Vh{#RP(UIzghM^i?cRn7T zGL!DlU0q$Fp@a@lK)11BWgw=fmn38F;1JgoN-a2;r)kjVCI;}#Iu4p=E0t1;IX^%5 z_V%u*;B%!~o2u6nKyE7s@p3>*T27mU#2;{Y*D7mDo`pKXd3 zS&s2N=r*;qAxECXI$P_@{~To!KlJu|yU_<+`^tR$yj|Fnv)QlJ z!R2SF{j^CsX`IOGMMQ*z8lnjcP|rl%;V>cqBzh2T@tgaq@iL^z^IZ{(&YAOCuZxNf0>sF(`I`T~54UiGuGQ zPu9BfPIFB#cOmD8s~mDO?hIm`5$CRSyIEpi%7hA4#?i z_z@n3iTx!~3W-e_^+n^jdegI`2~Bf8J{HaXOjrbS`vYJBe)k?VB#6Kuk^7r+54plI zv4kw~Ak=w}{t|nTes_}&DZ$@Hc@MZkQtuzX`hJqaf|kFeXo4t=3HQCTUez`STsbl| z(FU+>L(H`1zKq8Rhc}B#Bmmt~!oS52po$)l=(!|u&6}?IlEmzn6#0KhKu{HqNugQ_ zAq8Cch=dy~&jk54*(mmc+JPcI&egE=_GW2mB%wVXs2DVfe)&^ukCZ^qJSs2+6PR$l z-vOvHK&2616Yv4joYXbdnZ^v)B##v{F&D!@o-vo z9(;gOCt`4?dRJj_8748NQ$`rcRE z!?YIxjN22e^#N@Dly)@)3=fP3)Y?g5#Z1z`aEZ4hz*;|jeKl^O0D$5j(gE7~Xf)gP z%R_e~K}1{KMyfj=WsZnV+@WG*C;hMt{(YpS4Fy<>ipz@A)JdjaLvM9REs|c#; zd|?E;HZsLvxDl^`fQhqYpy_x1YkCei0NvKBuxvv#6l!?0fa&UJQ8Y^&24Jelm+N49 zlp+Qr%Xpa!W}(4!hFhq#2pB)KdJZv|4*T!v(1EoFJ$*y}pPCY=yXY@Ju&ycd_N@_r zcLD7`OCHcr<&&0HWiy0WiWT+JIKDx}X^B8`hX1*^+!VxB)J?cD89^NMvo50u+nEvh z0~P6(r!<#=Zx+y$TAbFX8>lpL${i?Je1QaC29rTyiB3!i+6pV!%80F;snSWd~!ZAUq-4MRa(I;)*Nx4$t&vO^YdN-ZAL zg+4wxxrT=~Iy#ys`sB%zyAQ3_R#!h}?fw1BQqSE?&UhKmJA=Kiaw$AV^Xv|Pv&$Tc z4-XHAGH%l$fPnBQ{>hy*h$!kQMD$OPthB0N$4rzXKtL*fTk)W_#2m8A)E$EbMCO3V z1?Bu*o`b_qpz`0rJt^?z{%Fy7xTq-~Viz6es z3dHda{-avifTSJ(7Q4oY6@9XQEsTcm| zVQf7rkY-~7h-PD5JkObL_&yu8}Mejm;O)I`33 zhS-9gx`+`Hav&Gv{cbu)zJl^U4gflnSx)bo(RF{z3qOQ~RhE}e*fdj*OAj+SwE^TJ z1W-GfrSeX$hBEVPnLfp7S<6i zddEZ_TUXrWq$iKMoZ>qq<&))ReonQU=(diyDN+3C~j}u|9$O=YQ4Yn)OXM9P5rMV3G<5fTkh4S31xmxZG(|*RC`f2l+nt z8kZjnWID%tUeXl=bUTx>8H`1kK$aPR&vyh6C}q{xX5KXWw{Q|Pgu ztM&8xYg{tOCer}GxV$_h=|D?LNeORx6EV@jhw)h;jh7AwYTK|yAM3Z|uA&^?K(0sv zNuUbaUjxa#c8wHesNK6n1cC-^yKi=18d$L}+npyiR zQoVLX_0{VX4oG8~fT0oj2;jitVq<%*JliXsZUFP2;+D-+^|416^xX3mi&Fa%pL4YO2_BK>^3dQiEnf znjurgEl|g40-=G6bWXB9vWos}Z@K>!+tT`alzki!R^>Ri+!MBC@l8&5&s{DzCgv!EI`L%dACY1*f z0W>Zx8w{7As|75earXu*^5kH%kOzGqo&+xe84bI(3`-gV@PB;rBR^S<%+uSg)Evbx z`Ie)6>Ah45;hy?CMEa~Z+ip8iVXPw#kmMzxGFL?Cp= zG;B+K=1Q$wkG?xtpra&$GaW$1l30*wz`REC3qPo7csN!I^Lh!?(JRAlfI9lb)KsB; z?5}L7@ce}woL6}ZPvjJE>UVO8EnhABeuzt;-FNX^$el8h7^BXI(he?M$t3=>-p}>w zE+7s#KRwsOljFi+mXqd6&4#M>2$c65ch9^<3#s2Nf|F(DCjnN3BU-O=1GHLL*m&_k zam(-%)f77S1tQ1$8CeF-Xvn7VjghQsc|8Wto_&7&?d`U{0eLYdYg`P>6`=yBB_+X& z1Jf)qgym`|9r-u{OTq>E8g|_l7p||nL-HSDYCsj}o9=r1w%nO8EC?uLP5#<)GI=1B z6t?`q`ZffAW^&O-I>*Th+S+5FIfGwF=)AiAq}oFu0}2+rpF`A%UBOn5 z0z3wC=h>cE9(Z4)!Pr#wbD>Giu0wd>3-N{j00)*`#yB@-&3YGg4f_22l3nF>abkjA z_#lA0(LROJJ(?F!>YvP!{-LL$k%vxg3!w5wI*Aaqx34Kot^`TO?{Rv#3P4xEZHCmh zzfsVi9I){SVk6od7f$DcQ&Yq@`l{!$M@twiagsb&|C9-Klq!6jD!B1=wkKIE-tTSo z+=e~p%|74-1R*w7gB7Rk&wL+hS(x60c*}!62FFyAOSXR;1b)WpfATYo|M4@Fzb0%L zO*Mfj7TcZtT?buRgA4iEl`Ch9jvt}a2wMu`YMdllT%(eBg&YF_f2p!YQX(=s$iTp$ zrku-G9m>s!BLIqGuvl{gA~4MWI{-SrPhZ7zxcXau_08MHyY>k#MhchtiF6DY*>VKN zg}JX8a7)*>R{>uL!lz1cD=H|%Bo4%Y5%{m-Gte@DW5b5H}WvN0U z9cYZy>MCJpvFUztz{R|cg9J#F21N{19DyR}viC+wxmjPUXzbp3G)=RNf=u%Vh+k>^ zb@Qu8q^09tx*hi?4RV&Zaa%YjTqYtA9)|hY;U4(C?d+pQ8W9O@Z!la|so8;&%_1N_ zP)Qyd8mcu5#hDq1XgrWcHe3b2G}u7{S7^7`6&sQYvV4#*LrcMawMo)xv?$Nbg6Ond zKE`s??J@&gxKM{r0cB2cT^aga_3d zRk=udCMe`X2th(KxZ7FX@1WPRdD^wP)XNC&BPT&g-)*7qUwT`0d;ZqET94jo2C5Z< ziWLRB*1+K4;M7#TVJn=D?n9OL4%pmNXmhxOJiyGQg|)P`c`p=$U3{trOCm)iZs^qk zUk1fxVNssA@SRo#4`tytZZ%ncD-aQDuxHpn=W!0D0^KYkt%@L&E{RerPNO}paQECz zn0oBIllE^cF9DsTAMksUl9KkEfouI*(a}XA$u3*fts%dy?m)>Y7c0QTBr>HIoMaTr zaOyeqjhGdPwRlE#_LTwBMYC#w)}$-(KCC>;eq6F2QY`F$jQ3UsFy3Aw=Y!3uH&vA| zOAehXo)Gx~kVqkjVSqyT%e;PRCYbXd?`tbOog5I?%TOTq&{H32Lz>|ZL`n@9B>fZW zi^7WI0LO6)9ESpy06YTAunLOnK(74+dr4Xv2eXU=Vj6FtNw}c!3KPj70Fxm-CH4BR z$!S0$3jUAa5jRX+K+zr!g;kW}$r}VgP2#moPS)0n^2VyPGKYymFJy3i@+spB!wm0zogdE_(%ZnDQdZu_? zk8$3Q?a8WBe}??zE2d1StqwF=x-1nsQCyZBg`o3+u^1LdG8nqKxpjYjqV=s$idD;) z%W*p=J^kkiV+2Ml9fUNB#LzbE?Ar)5?_|Q>O0!f|S9gxt-Mw=P?zqKNA=R5pXb^tv zwEQDo=EL^C#YZb0h0IjSZZK57aa;Ui$Pc9cu4zq@3a~W!^k6-y{xaXF*LGVd(}`0a-J!%pvK_$YLDh`@QDUVn1l{FH5Oy1UkLWF|6Zb$2}yR2_G2#`FS5(n<=_ zYE&FOQl>A16EN!txl=)Xg=^A**YIS(hZA5e_m;&_D4Eec;QyMQfi&GfYHb+{w%vdL zR&JOGod3Nw5ftjxANy-y02CL8Sm*M7@O_?&rU7B2JS5}4ci|CS)>WW!SRVy+L-OnI zcu91Q5|Rki=%8v9?Bpg)=Julrz>*CVc~XpwuG1mQM~Z+?Y`Fynv{ntQJCOz{&Oh^8kD>Ao~#*a8>XVpg;mey!RD7JCVrqM@qmOTPFZ% zKR@;HO?&t089;zQ2mx+q)h}R(JrkM0zQ02CE0%LnwDR|7nLueHAx|h%8eW0HNSZ2< z0}7D4DjR?IZ_j~u3xxkSjC^h)kvDNtAA#tG2Jm2Zg!`Xs{%~lUlx=u-(daw4k^om4 z?gK%o*yF<^=-jG6zDfsKAWooYf(ZIuL@JQvAE>Oxe2fi~AraC6sVNC02kfRV@(848 zs}hi;wGZIo)*0w~z)vawyC?!o1~A+%5Q4))Y0JTf=a89_d6L2-I6Ped+MppGuInwe zNlTmh@<3S)RA_Wr%XE=M_suJ8f{ii(ro^TRbSPo44;Vav6OoX)!cN;K=#ZSH2p}4$ zIc!4Q+FMn!++VV#ZOCZ2) zNPt8`{@N1gJCc%D0V57zii(r>eJq$w6G<#c4|qg?K&}+cu)X(fl8LDRGQ15PEVg+A zWB9*q>}jA%?4mWtdLr>}0re5kl`6pvQS|us)o1{`P-XbQk2xGIEln4@46XwMLjfmC z8vd58FxWEqx3VH&I~Pd#p>M&!?nh|EF_4fhnB&y;aak_VUSD571%Q@yc;f5vuf>Ih zg&=PNT|dkkw?P1W8KSIkY~2AyFMV&L_CVL7zokXm)wP;WP(T2rrj|mN)oG}7o!y8q zf-P@u4xJiw#`EejIjU4^wC_Yfkq%gF<#udL0WZ;$J!F}=0u|a|QC`U=<;86gT`L%OEoX&Lr79)+t?Yb^bx2E;x%>Hgl7bc%Q049gBiC`q za&SyilwWS4?%u+~$+eQ~tHtc{pf6D)*+E_k@7lHU7~X=Mz!qPTS3seiI`=!%m0Nie zAgu5bQN_`Fd)Ry?)?B{&aI|*XX$Ef9Lgh*R`i>z%H1Y#gzDYhSa9QdD(u^b2l zS6(Q0N}!d0Dl7s=^9I1+#$8q0tKpbAjec-rOY9GhjB8InC)l?MHyRX*0YreJC4>a&!b%LtHA5AtgrBF>o3~gWtGf>jcC0ESt!+Zy@#Q1jMdIHQX~NgUdP7P4x7U)R zUxGV+(P4XV96XTA@M(zP0O5pu;yH+MF2IlzlPFmk6pC!k9<&7gyP-X>q2=OC6R;tu z7~!z}sMzi^soKlS8<*?QWwPFBYs<)YotS%}othr&5uV3myLw}4qCX~X7$Y~1@!AoJ z);Qo6TTb|y`@BlbXa45`b~`)UjhV+4I&LeyH;$-bt(a0SGMt}uioP1hBi=gc!6x~7 za*GO|D3^*8Htu1F58QqIh?$t6t8UJu$7%}`?Xwq((X%^@@R<^)qMJBB859FGj_2kX z&k~N}N>8I6JV=~GqZ1v8eu;i5!rh^jOR8~O;oRi>Yq%G!>L z#$E%3?iMGuTKnmcy++Bd1PZ<1_S}^fLJ+UjGYvYKb zB@2Nu!#N60@9qiXw~KLnt#=shWHKuC>n*sXe+IIimb*}s4DW8wsTd3Qd?pxEdm(JN zYNGBy$hkp~eI-B#=#2o>sncETvGnI9c|?~Uek}m-^2^4b#e1KRgB9YAq;}siw0pR# zSI{(Rw@dNS{(C3)j2LPta`IR@Cpvqtu8~fw?!HB^tGjD_k66kCmq4?(IE?}5@^zde zU^P8q66*8H!gH3!NF}PBG63h|QiyBCdj!I#8_V^hRtvOt(fL&?je7B}`!{<@Y_k|p zO-PfM*P*xxv}x%bUex6;>cvTIEockRv5}V?CV=>wiVm}{QKc@tP~heex=6e5`f_OFulm39DI=El(sNvF^ho_kAPLMM1^h7asj24PUV*V>+W zaB;Hhc~V4K{y8yg@!ksadnD2@7liukp0C((cb+2TpixyvV#3qB=S&$MrfY15^G8Gx z#i9spTy9c1&i*j%N!5m<0RA%N&wk+IDh<@{V*7@i7-)X~c9@sPfLBQxrN;6V=XgaR z6vBUsAT73kQZYABr=Ul5R!5mJ`6Bm#2#>~gkTv5m!AQMly*+xS7i4Dr{q*t2Acsq} z6%Z73b#*Q6qJTOml5n|69?eHu~OCkXdo#m@fx@}6N_cJmY$xTKY#v!S6=x)Avf95 z)a3gqe7M#q$X`-hEe71VH&Ej=(>njSZ8ygLFi&T-r(cv&fyTA?U`mJY_`bDL#HUJ>Q*`9q`%E?fz=I z#Y~#x4ze?iB2*!#Fj0`enTB)=wjwpqS#0c~4d@)tiFSAFo;XsRWjTbTQxRkvkspDh zNXCEUpa%>oL)YW{JLeL~hM#S~#dmIk?e6yx;MZ)piJ;!o8kBjl^1Uye;Lb{Cy>3VN%8&;YJ z+Pj5!GCPEg_XSEZnj}h4NhrpPpxD0lo7yHV&uk*FC{LME%-euZ!Gm#rb3Z=Zz`@aE z#JI3gpW(Y5auL8nLf-?PaB-Ex(v!{-g!xZCJuyZd&qsN9rt57>auwCUJ><)qB z6u_k)Mi=~Kl^4BncaOf|{tvVomg#eljM33!^w}nIX?L}!ex1^Lsl@r=dP$O9TJCaF z-~P^SKsD#%?LbQZ&u*&2F-ZaTC^#IBfNJZCxXO^rI^Xe_z|mim6#j8Y1cQl%RRNBe zaG0C=OjeqI=y`PaRbRY$^7F|6*Yy{aFy^9NN&(zHMvOIE?b`fe@Z==8>J_8G@$DDG zNdu`baiT~>vHc$X*Cg{HLvrPWo1od;I~Il_m*f1ze>78P3CxMICqnP6DP`8a!jCUZAe;BG*h55zf5sqULA0p60_9&|l^g3_}dTSUg>z?%(kjB(GExsG7wR22h>P29q z(nMe00x1n09kCvp`^x;H=XQ4AHwvTdJwRV-u9&40JbC$o0pUiIJfcqH7--YcY!!jH zn27MqUwSzF;j2+FJBK)sIof0U9)0Hdj1Yj7^*<$g7u~xBdg|$}C3K*pTpTyT7zj!d zzllU{=p*TRdv`jmi!(c?@SCQ4qHQ}~9}xptAT0Mq5j#V=1L`op!8y-Lg4SNdEpJnU z9vABNIwQN#~O#E0n#Jz^&)(tMnc#fnd^`603O51l*9UM7CPIwU7k2GpF2I0vF(D!2 z{A9ZO*%L?8$I@cZ*`kHtB?lLg!{NxM@gO`kMsE4lA8&snP!WBz`|YjS4(J#^?o+#R z5#wPBYkwZ=*}vW)eCH*FnG$eAKV=L6`K#Wgi<}n|;UA}jdG7C|2I)N(esZxD$$vi1 z-yIz`*AbJn+FI@6LB}b4fXxw2CoRpUVRj3L`0X);Nx3xd&Oj^e>7gG%4C74A<^eT_F$DjinNd^&`4N@UqcVKO&(sYMZ$xS>y+HvWV$ zrcw0~2~KZ?I)Yn~!Kv|`ilOpx+u5CZK{w@XB3Q?7R&r_K$!$ix18ZsFy)E&R9lKg* zYgysLH2E*zp0gyVWmuA&fuu+ifynN;9yOTFGu;!fn%+vV(COSL?BZuQTj9XUfQ-6& zy7-0mT$ABNgr@Y5hGjyW5yq{xU5Z;|7ka|Chq~>^?iEsXIIvRa##xtf=+#(T8F@?E z?Tvm{IjPWbIhIjT`bebfG0D_;*hY`sE>1VX2hXuSOYQh1)o|^kX2Y5lJcagYLQA^G z`8(gdU1uG&(oKdn->ff7d;KO>emOu_*L8k=QC@yY>UONn0BGv^BAYaX_f_M56 z|4yX-EUD@}D%L`+tqKQr%ZHVD^=1+%(ok;m2YUux2nQBkiR=zzxzRj13; zpb<4T-ukhF^7J;SOy%%ABMlv@UC60OP%|wzU(bzEB@bjV0p?jtoAlWEb zE9{^go6(^6;K4YV%&51BKkTvRaS9{VwWPiY&wEh?RINMdyXDWy&ojv71*6$R&uZ>p z#-{OpPc`3pJU?!XOB^P;Y9jd_3Dk`ryYoRun9cf_V4LhPEC6m)Ju_Bw8l1|xaTeQ4 zRD0^CxwzS_w;MLQ8XDb5EQ0>glgfcxBK2orMDAK9JypCl)2I?ZmOS_1R87h{<5hk& z@UtdjkZR&6`9!$m(%=uHXujXI2;-sg2ORs;I3U53qhp5O1GADI_R7hGXFqcKtj;2G z&OYScn(Ofu!wzf1fz5rPg^7e#5TH<6pUd#>B|coj!Xg!f{w@IM7=4@cChLTsg}uK! zYtMk?EKT;J0Cunv+V<;i=eM_Wo}XNnrs9tRbOp0whPrD1{*t_#?DhjZEAjR2r^eO4 zr?`LQKdGzdzOFyfkPHi;c^EP1+h`Hi=d$xVyS`?s=G0AASm>0Y_D4nDjRZpQDwWFq zI-}^9S2(3YQ;K1ComWT6T*d=ELUQl6+}%KQU4O7QOEx6mDGZRsTWXl^AJZ4|{&&xT1eLM9()9T3Hn z_k?kw`a$A+`5LF^w)b4po8egZI=TPO)p$D-^jxMjRo+bIPJMl?5|)CfOGy4g3ZBp1 z8}}k1&XPs^T)eARe$OpqZ+1XkNVUq8matHdD8_yfEKbzT7e)S7Pnuz2C*o39V{31&C4wXD5w#V+bN+d ztR@8R-$!a_gD#w+S3W)iplqu|5(v%$7A(reXct}2ts!Cl4fIXJMyj3OSXiu9SdM`w zUw_xfVqmc1>F@mQg*rP5p04K>V=t^N-}7Om-8c{PzgXyVI^R`|`Rrvr;R*Ztb3$HFL6ru@bGR&UQ84ZtdY3Q_bC!+zcoA ztF4Q9v7@|4pY=%q&N;@YaxJl#8(GgZ&V|n>ek|-VUc5Hz6_ek*sJ1KA z`E9SvpSzHl*l&MC)|pActEa)YIb#u{_w|^}E{flCgR1%G%Pv)Qg;pPNeI)~;?b^eE zR1-WG5c0mMThUs+_6D0-@MTH)@;ICLA@olCip_jm6u1Zsm1{ZF^Yl2EJ4Ti+zIc2# zqqmg9@+ouRo6Q#Q!-ut}w}ic3wg{VPXNG{xZ^*4sQL7T=hVT|aWnhyAuq7GgHDY#f<5qleaG6k<3(5FV^9vl2~x}HZ5)Cg9xoU{`+Uy9@O8ThXFW^)@QDu2 z_24%L7B}g|d|zU#UX`Zv$6$4yP8}|OPaCuEeHq_&IXSh2qQ0$-#yEB+CDt37$0UWt zRfMeC)EsAC+4-6z8|f{=;H}!9)OMP~-z5{MzYH$r6xqb5?KK~-vZ#TEz4r*ig@W`= z*NZq#4y!6SL9gtAI)&Bwo=-hx$;g`8)PA)mA#ve}$JV-(n#ZPH@ES>9+0sa*=R)VE zOIYO%Gkn8HFS}`b8|V&?$7yZ&ey(p?mrmifbK2(WMapu$o}kI;w^lMw)#JtbzLBHQ zhhzu*H?guGGp5t+2As>a&jo#yXE(7rVLDa{3QP=vTPvZt1C;M|a&BBK4<8=osTbaWPuobXet9^zWyRPfJXVVRoWEjyBx_j830Y-fA{y}W3RS+CvzM~2PN@d)u@Zk6os>%` zrI<5aOHEC^^YWi2(d|N-ddta$mwc@4&wrX;H#w_bNx;mxZ|W)j=3a1{`ME@R-57cE z1idIP<0mBWpshUP96Zl-ppusE$l#G|OXhAHulMw9bbq8b-4ae;qI6||EB%wPN_U!> znVC|Z^P>;&V^>oP*Q0f2RSSM$H9b5v>d`!r8+4!UWQ1nc)jh4DZl(xbHV{!x7L}ps zWQbu|*)%#k-tQxGZHSB_xxfC&zt<#;HR*|f?o)%I5fQPE)%%PngZ!@BqXW{N#=p5o z({U#2=`$tY`|_+2x(4N^PCXyA`67HnCumIUd!kII?vs+bX`#fzFW0)4%3pG;zyycE zGXU*3wfkLyo6moQh+iK-nwzNTu>6l;5Z< zHz#Ycs-uplMYwd7aa1`Z#ws{U=``EwXeI=+f9y$4Ax%Gb>=18yUS?EgRw#`sShzau zaHo5w0NkxzX)^aQ{H^OAPcp)tl14_D`oxMbT0>f--t?zh4BoM;qPzSvdDev=FP2lK z!d5s%uk!c}Ws95q_L_L4C$D0=oZgW*7|PmM79OUxFs4<0-SFWQVdp(!9mm*4NvX&# z=A-IAT)|bgSqG0zFAPIcE#uaXCo<-Ni^(d!~6xoc}{7dG2XPFg9%(C@e& zs@!zNjpTjBLDgaQygOxxpP>+*qLF1-(%XAAR~2uhvr{r4qp|T6CEC~TL6W zM%?B9)TcL^ufxSpNe%2zhnyT!y;&!yF=RE^XFH$wPD|}q*E1fOn_G*~Sl2!ejNLZu zqTcPwedG(LlH_)DGe zSO@Fsl&EB-q5Ho>+atpdSPS`bqQl{LHS@L7pM`~P?tCUSH2+O&J;dze`)EV2v{g+O8-3cy@8GLrT^!D|DXIJA_&7kaU6w0Jj#Kf$_qDb0CaLcZKxybN^CCA{=W43ZjeV7 zb1i$uik$wQwA+4G7<{t%kI8L@42sVnUim?b^Yh9Vb&^#HYk?@&fP$gtq4*dl>Ed*8U)J1xELj-*nWQkG<1DA$ zyHkh$e5$`tksPe4RG;?Nb$LzuJo(Ow<>P!ul`1-)kB(LEsH(3g;F$m2nr^Hpk{M>v z{a7Z)pgI}ZOZ=36vZKUt*cbFe;b%%Q{^E&5;xH;EQ+h2BFN{uE)EL9#v11+BC3yGn z)J{qyYFme+d01aM9?6vwpZqY;HXihK`t`+G!N_^@Mq>@7`p=Lj8LMIegxPw6lo9h$ zA23(EW5K)b``+mzH;6G9X?8-YM@6q1@hl3umR(5moo?nR8@!%wfH8k0xpB3&lc*9! zcIC>IsGH`Uowr}6?6Dx`dzzP$qAMZ|{zo(C|IgGO$8p^1=DJ<3xNdZba3|q5n#)yl zL%z+#WHzC=6wT#k$)-r59=^v~HY@pl%|_W5kNL7)x7O{8h*{FKb&y zbpL~Uf8J-0bI#*&KIeVj@AG`V#wK>PReh8frEi2TH+3oP3GkkYcer#L)$&*8ldeIP zXJT~E=}hXqXCZ07uEhtjX(;`nQg`_uPKO-Xa%o3tF$>QNvU^0{W?{jRn%0HktE_7p z@)tw-tw1SyJxZq;lMFUg8%`*f zpQzDWU3T^v@9{3#k1P15pzm0QtTQt(Z-7bJJ*)&2=Mwu1+kockKu63*%Fb7}bTd!( zruvErYZE4%n&;w=fr0~~Y5apOLTFvl0Wejl&!R*AkkMrZsN5dm<^XWoamWPpbYWhP zzlC4}Uo=Sj*zWKFr-FR~#x1ZLNet-kHo%nkEu(CL7MT?204z7u66Hk+QBdtm7)8W_3UK zn465H0q~7)_@-0EHHL$?6~&?uSQUVS!v|CNRM|jheqqABXlmR^$szb*0)A5jB783= zk~8fCuAQVi3!v#GeXsXx1-}{A_Ilw`dIUgyGnh5wXP^G-dUZ8OJj0ZYC>whdA5phh zAYdtL@z6=3caCl11I?O$F}p=McK=SLd5flBjFll zvof^zZ!1lTjF5zcFKUPoxrw0+Xx49Q4b zPi}urZt7FZKBbD`?B@Xj!YDop5GU5$?K}MG_Z}k$9(#d}c~h6qWS@t&do(!MD(RdnR$7-Qfcqq^r~%La&p^)i1&$azWiwPI{q7bnwlXD z@#mD$^^A-_QMhwcS5I+nO@Mr>eXvq&wl=2LobzAoh*ZA{o;mYkftfA8jW?U~q&@9E zYGr8|D>#+NK6q9MCdMxZy4EGrty?+N+HjF-sI^0j3|YlaNqT<}PsXY5*LR)S#(B)} z$($`=VK}=~(O>^$M-=YQn!hRdW$aYdAoygR)qq~CL`kKwSBFt;;Ye<<#5_7I2DsL; zK~uLow~{57h0@E{1FKO9-nj1`Qgo9lzs7|JN;}!fE>@s}KA*o*=Sy?kAoNmsw#AUtLuBB)ZSIFN5xdg z%X&9bT>9tv4Y=Y@pV-tDIs7zvckbbAYK7Z|4`sGaS5LIO9JiXLSQ`gE=9bacXyRK& z60g<9la}Jw)X5T{+6N3PIO&@AgX=@c-7zUgKiM1XlA&aVU=AKXmN7EMy|ejZ`)kb* zBO>O_iOl%Oc@`BmbG74a(*uUMdOQ=AK0((gF$M zF0>2p%v@we#byJVsZbFK4Ce@FcE0n_>7~@c8}`&#UgIDiuW-LMwb5RV4d@uu#VQ2Ej8mLM8;ew6##O8xfn zd7jVL_g{F|axHG*-ZgvXnrmjyz9ZFCWU=p(-$g<~!j_kldWM9AJcon?5kiN6-!Od~ z7eYd!L6Vmef9`?2n~mv0G~u~FG(4=+$Z*KnJ!wc=fy`XP5<%$T9Dm@}ILdmxa}q`Eec~Gf)I`rH6Bpnj|30 zH-_MO&9)pb4RvHqWI3-7Wc2tCB%SR{#oOPWX5CJ0WM)a;e!TO&m*FH0-tNUQ4{mfE4;mN(x#u_3tUu$%Y$w-;$V_NJ~4 z1NN`Gfj;G=+a^MX@^<$2X~^taRi>vqQ*qeD5>q21+!ljbUC|^$RKWaC)vAyZ;@m}joiQD3? z4}LH)HpcVhi9kVriAFguK_|&GY3b01qnfRO$6?qMg(N=G$bMc6(VsqjvNRnjl*b_* zU=M-?_{HcBa{FE7U)DKqMB!_e7Ax@^zBDm0@$|eH$aq|-z3|006Gf!H4)elmvd-Cp zT;{EaH*c(VsRZN@-JyvxKF0atQ1nNY&r6}_Rk32NmD767K6S70<2j%E&DFamC3~EN zn;pfFJ+2Uz_1t_KXgz_pu_1S4~JViDk5PdF13KP{NaK9~9#- zz9#A@O!`u4{%RJEA8I72qmwxnn&Br*p`Lcn=?kj8Z{aBmYb7u-0B*`KuT%(!r{1`H#Z7;TN(=Fn6R7g zZz}K07jRS;kH4wMLn$nrjBmH5I#Rw*=AEehAYQpa2Kkkf99H+!Tv*(mbzoc7#N7h= zZ?c{R*dydOZ(_J0eD9PQYYey~wI#{J9ulj&7U%QIwLB6$dopHP2+c3ix-bp!*G*wLo<@LMLIiFARD5o;SmY|A!g z_gxY*)58y+wzsY38$EkxYwdHWxS3Yb8cGu@H#=xe_Pdib+l=a%1x8MmI>IZXyGU@1 zQr=JW5wmHB_>|@rao)5G*Svg8$ZSf)YWAHYHD8n0nrOAx`gymV?y=Esd8ip9OO4um z6n4DW4Sydd51|jOJtaxZy~ksmCi!J$gZ%ECL}E5EEsxYlXJcL@$dFYM3Y*EeJaR<8 zgGA{+{6;r6VY~aOVqDp`VVpm7FjiU_;YL$S&y>Vyf;FRJNf8D#v#Z2}81o?2X?P9$yOqkh1n zq5haB%f=KV`JuE_immP=(~nqLF1M9~wS;?NNJjhzqzt?F48p7O7+lnHjnMOBwC~4u z800lXsW`_(Zf=jfK`(h7lj=IE8ADHm!*#9KZ9G!Ng=V{rfm9L`KDKCS^K^h;lkIDW z33SuI4a+q@t^&QCsag9fB~)oZ6<_=5P;s^8vjnNRMrWne?(QhLm0uTwhd*y{u}3I) zO~3KmjI*-QO1x{3aU53KAcWv4KDT7&BqvE{GI3G19tckrAAdTLp7CtDQ#|BJ5i!9% z-roy^_o78js}$~82>cjzS?d{nbnt zEwAT9oTm{pu7_-p+nVjFdtde%9AIs<`UizQmk$joG92D%iFH$zfSa`evUdv*%3K-h53GJ`BgwNbS^ZF40?hCgq97mp zXkeadNl3h^(ZZ?Xk#&4?YxH3YXK4JpgU4N)_s6?UnAGl9mHj}b{P%wS44 zUCn`)4U`XvDzpo~k~8I>?94v0^n$|mRpK*6YATGo(DbFD)B>(*y+PDibT(*;ozxlwc?M_iQt?;Y-p9W+M_Y^n|>Cvdo>kf(8S_K0332Lahvy&Q- zI%M^V?Ha}SoZ*^_ik~eXy(m$;BT4bt9E+djD!qktd5TQL{o#xFiPBXR^7IbgDd~Rh z`;j(FAFc?<9rUmT%N=kcW<)yLn0C|djL_h9D|=c!H%SS=3%Dzpc{~FtAdY8y>CY4| z3kz^aXIMl+{>?=hu(QNo@iFM+_X>&CENXPiJu_U@N6PLZ$Yp7IV$yZIc=@a?K3#a% z*O@jAlVTX>slWtl)%ybzkuW@71H_@q=6|CMVOZ3GjF~N{u%XF-=cnK~UWM`-6EeX8 z8&DGa_o*oK;Q0Y~mSrB?HUDXHzwXJs_}%okh6A5OcCw+Rc)9F4d7a#L3<1o#ub%rC z1g%(+zp|rzYaI$G;B+yoxt?JCmx!c=lGnt#N*S+kOQpoFLB=#f+8L%Kqq$8AV1s_>|6*uLUx@)j#(yDO ztqLVS48Wt2r&^c6alzgfX{#bRlt`3e!4v=RMQSzRYRO8{bDg5f25Z!42iHtZV5&xV(0&i~c>bNqt268F!6pBOX{!1Tz`1PS`jg zpr7N??;|KESPf-tbiBKaWrRyCM;U(4+HTfEt)X^AcZ*_>xU#gL+uJlf9tBW+6k3k) z&Iazok-Muj)Uww-Y4mi6**wSK|Jd)bP&CcFFVLMidoW04+Bf5M{D2VA1)e4(?l;## zyNL(efr94>X5?CmO*)ck`*2=EU^W3x77)N9l=6p}Cq`xC3Bi6}Ck3y1k-wXRVhwTs zpSccVJjbznIpFX-?Xb$I4j%5_WwOLdg{4QovU?~`(m#%7Zg#*biP#BthQ$y5WPLN! z(|7NFH@*tJyTW!D8yhP(F)=~ED*@w)LWlAdB<7(k(f0B>~u?H-4r3M-61K>z}wYSbk=KtB+LfO^XBksB3YLvqYMr~ki6Dru) z007Ye06M{#-*y0HaXbvLgf+;839?s|NlCDnuVIwVIX8=f{L`(y;SFg zYkyhQ8uwUo{o~#o6(9Q_4&F1hB{dHI{w;d8mG==Z0vSadO|;aulr(6tva-eboTw-7 zztvun2;S;(SStL^0A9MpFf1y~j~N$g@M`&d@!@)Ia+uC_3McASFxZn5=wdVHl*dtw zK+kg;jfSx`-9(EPu;+0E0jg7D3!8VHLjxRWa6J)zM5LNfuOn2oYXY3FwLbsn7wU{OCi*p&Q1?W0;G6yA2ZIxE0Y9R{8xf_JZ9J= zcWkFl$DMNf!*T9g+t>u-^&5`4Kh@Wdll&5&fSm+P3JHD3_j=mP^LM8aJ6aNo9tahF zftXHrFmxTL>v<$>CUmbBHHYruz)+^}{O(UP+3SApZwGpd1UV`KJAznWbvhs*tHl9*NmtRqq1&xm-R_Igz&x!JI3$2q{5ad=?;y1we z!d;$4f5K(8QfEbg6D4G5*I(ow`GLvUtf<$~Dh4|>e6)UcY^>U_4H8;qIh>!uqQT3< zLnFIC{=GoyPsq*&J*94N8@8YB%=^NV{zHaz<8g_{6hD~89)`==vF~!pjQ1t#+mp*f zu|Z1xmkHjP1lkF~Km9Ovm&E_*Xa*8-KV6?kR_AqoOq+U|ttiSQC9wV(^bdq9r#ixm zok$|vRWmBI(*V^B9|vj}w5;p8-DdZ5J8f5evp~uO=I%g&j4Se#Zu3LGolvMHE@+(% z6~0#{7p|=L@*ZNE=LrGf?L-fKx2y6EgMjLU|M~NldK%Fk1T2d(EDCIJuLIcXAWW15 zz=sE!zpGLBS?0gyu>tdSx=+fm5?33!g@;JDiE^uxK{q1M(`pZ<^(uUF*D1MRE3g~` z{O>1K%wQ<5|Ax|l;+1VK9!7IBlnZY$k3UQEb)|FS1cNHHIxj_yfkpf?B@O6Uvo!9S|W8g|a;RJ!U4|1_ethNB204B*;M8t}0iN!U==Q!O%UpbQ1^fyJ|h< z6bJZ64bjKaC_FCij|##7Q@O=EFox3=ss6Wy@BsJh{&bY>Dt{HJxR5{1d2|z1sQ+tD z2=Uses{LC^%ur$SGP(bg5^e5JGyA`%M2+sjr<*wPo}J|`6R_DUSZl*l)N9fIdPIY4 zcOK~eJ%_*m_m01|0T(KKEfuK*+z2i+T@TPe8mRF7GCAhN|3n~+tQ!882sSW01pgXO zShW>x&};zNb?O&09onDfjZ|fF8~ULbTB&=hW&EtE8fCf<1-dVJvUnnW5dqPgR+5}a zdfZEq)2SJASS z%73~_ik$QJ`#tFAtbgmOOl#Udx@y8q_l*$&F*leRg^v~gHBXPoe;}rmb3xYtO$$1`6b|d#_6s=KX;p@REzm ze@J72etz=LQY}=Z$Q2Mvy#he6k@nwGQIT81{t^uz`uUT;mf9%Zy713ZDVXU#F(8&I zilb4euJqqh1He*uTK<+b5;FDQc8-R=={4p5w^_MCZa^3R!?s$&1NXmh1)d!s$mK)! zm)+ViFYf;}fP)b7g4k0!LO?(-BjqjuR)qxYcMxq!{pSJ7m=-Vpe*X#b!rv1r!GLzw z#Q=#l{FW+o1QL^=K?f2m`wwf+2)+xY`BMHhfjvgNkiR?RK-S~9GT;9C&`2Z4vK0;kJ9^ieL#o~*H% zQ0ANCu^xNC;I==1J%84GHYl_oGc^se=x;JJ*S=p}{IapJNz{is|fkVk3K-_ZHv=6e+7hjokgnfTkyBR9L;`p%a==#*Q0bim) zxo~xLeb+tc=7l2ptMkji&)6#6rN8B_K zAtE{yLIhBA4E6SCxhtAk`OWi}FPZGt-n@CkWzubl*4=ru)^Fy1{rluVyHvkv+G9-N z_s*hu5ko(M2?Vt66=;~=`-dEz9!ER^)JU>IJZ<if_JbQg={}9Z;TmV~SwP4EQ$6KSxQQT3H#5N0Z zxx{Oe=phIXz=&7vAT8)X;c{U@eMx#nT2xdtz4`i~`ayc+9yeeLFo+d^0T~$Z^>h?2 zWU7g+byvwbS5ie}XXa{YL|h0;w1tnZ4J^zwnB5+nHUL@*GW_3VGQiAf0tYal2ayDU z<{i$A@QAbvXvH1GiZ}_-4dHb6o3VecRhsnJ%{SH+6ciK|N=#@6=zQ5-BGs~Cg(D6Dfe%|x6Pc-Iah%qWuXwrev z;Qq%jhyiDD`T?(MziR#Z{5OiZl1gn>dv!3>N7hiaA682DYLpo#St|GD-M+!iaJ=8x z_;}cdw`X(i)H!q*WAYtsNw6;%n_dEHpigUfON)W~Gi92j*) zD+M0Q_+R907p`4v&=j0hVW2%Heac7$vdCIrolU)BwqEip-L?IMHQ6~3wvbcYG zGuL}XWzGz++W=@OA^H&^qr}7O%hghKbNN4#lcVWfin@pYujCTEA>s%?-5oX)HmS0m z;nh7Gb^B*4c+jB&x{1qG;%Rxkx5uit$7AZM_xT-=cFQy?Ap+CLfAScH*+|*_hGxO+ zEeGwGxFW}R`jNKxMHK(>tAPOkJ86JD?7bZ4yUd##98|H*{-gC3DlqC69CJ9`9ne%Y zTaZZ6y!Y^MrDvktU}1h<)cf?+Zooe;f;7&Q6L$iotY>+Q}+LWA>ml4@53 z*OY4&1c-|EST`}KYNIX*ob%a4|4}d)`M-qw8CUv}2+5zf@QsX`6T2u#c^BzB1i{c-_J!8yOiJ7q{dtD&U(}=lQL!CsbdVm(lA{$31(jrPeP;e2pCEtt>gwFYXSk7DED$z zTT4vJk*ePjA!_ntW3xAZvScpo6pm9B`tyL_eNAR=IWFR`YEVpNT{_aTLPP@z>A=h~ z&L?yKoLN5U1KNTxAK1Pwjt88M(L9)1a>n?NvW(EQI!H5Fl6^{}_B|d}_5VHUaZ+>0 zB;e9|xvKQKP7(AK5Kxwe0+=Um~FU$&W`25mQ12bBzj1 zPGnyRTZ@as5tSf_O+k=DVzBmx);~UOZ(C1RJ1j|5DmfdI_bGw(-{Kp>^V6N*B>IyjruD>- z5q_TY$N-_aL>7H#PBP@XC3I||!l(ndMv&uM{$VT%oGJ%m_5)Vmrxhl{aiXq30ejQ@hMgEeXbn73 zvctiWucUzfEPpe9ceCcP+r724U;SS#z`#Pa6`?LM7cLiwD`_mD$Ae~I=arjm>lp`(8gRkcs6!sRp*Xy7{6R1meqXf~$tHtcvF6*{-{?72sO zDjOn2*xD6X5-}~&7;+NWWfOP-Hj$myVLQ#cGil0wPT13_NU_u@m<$;B^0NPV2ep2M zp*Z?33w^bnARCI0|w6yMFu3X|VygrsefU_^2HT#{?jvIkz#YC9 zpa>>}k}1Uq@CxsTQzSd8O6%xQ04HUK?rF3WV1=*~PM|1sgm|N6z?YHedms;su8GzG ztiG-6p&FHyR?2ovLufDzvA`Tu-qN%~smL(H?zQHqsJqR&kGNc@7y|SN%-FFAIzF)KWJmyZdT7I%T5G=$Cwmj~ z>6fP%J@z{gF}y~&=O}1ojLSvw`?nIKn%vBr0SM)lgCT${UoN3y$uGv<|8fS7EqQs~ z&DTtOW^Re|$1Bi)97daiG=>TWaPKNuCmfQAhl|wgNKBhah#nw{KwzaxzA}C>L@Y0R z6-y_u?M43JbZ!7JRs5G^ zTUE!j+pj9ch!Kfi)<3^I0lMhZw{X4Z9=osSd(ejFnH<)IVaP>*j`Qd56TrV+w&D5( zXi8^iXRD*t5;Ci(y=EJ+0B_tw7EEJPsl~3JouBvXx$nNX<`?lO+%m@lyF~cJVi@vg z7zm;p(SG1E3LiHXL!YiCgek3fiv1+6U;vpeJoL@%e#BPR+iG&q4>Ci%{f?4zLZZo3 zkp|Pbg8gC&XSD*LK{ANt4|xEhvEXHu3IZdVz+-fu!$^@TQ@@?hwcp<8XXw)-7n^l> zk)3f+-0uTs+2vgb35xFk?^Z}4TC0t4lU=0Y6QigauBgNSAdKdrk9~JY^)RekqZ8r( zfJG^6z+~!K@kJf;%IoSp7n*%X4gpWR*fCIWZf$2+^oDp1$Ahj0TxII7dgHGFZHcdw)%Do;*bU+9iaNUkY^6adA$GKS;~RH&g=Dn zQ2EvcD6r8V7d60!JHMaz!;ui|_czb8Y}8a{WNJdt)pO`&g})Es&FO3#TWvwgvFFLZ(2c&aUKJLQ6Q9}3?LNQ zSn^OEZifx-c5m(^U4>(f**)6Nz^MZL6Nya-L< zjtO|skrF61fEm!A#z7*9KZE`vp&RoD+#Ut7Q^L z#V}1Ifw$ad3c@yH@qs7YA(~Oe`;as7{*w-m9~p=siwJ>@UzN;OdjcjP3U{A&8$4OI z5WCcpl?~q`CrmJevsTa&QemyoB0%d(3J@D}cs|d&Cj}?mjd&?3DQY@!?)AZJmgeZx zRN~q6ypqLq-lZwR*>>7+msqR~nPu$<;9My9X^$R{;W21BTR zg782W5t(^T<%l{?x}w38?;>wIJ*w7B1q09mKCha|&toJ`rGz-@jlhWq@slPqF!^yU zMs|5Um4R?sAfi%$Pjn@PNht;F%=Dcr07}NeQ9&@k0OV+$sCr5gY+SG%;Iw!t>REeA zE~l#iG9?wr>B|Koy0Zy4($q}gD{q#NN5NVNm!hx-cJosV0&Tnvlykhka> z(6dksJD0zY$6Cb@{Esi31vjGi3I|TA{(JEQVMjLTw1nG5%W~e!O;+n{%0XY;;~X1r zt=;OZE}7ePD*Km2iCpVr=qzO!`uysBA-6>iOPq!yw;xlxcL*)8u6Nw^aMJv+(p2%j z(IeLqI{YGJEN?mmJ3@ip?#^{GR(^Chs-MbZU>CMz@7NMqjo8Q^N{`Vj9~0#b3b==$}C^N;tnYw;}w=f1zxsf^fx-a3nC z`z8e#pu>4}huebD@IT0LE;t^gCCU}<1ir`yir$5X$Zcd}6x| zUTmT&z2WG=Nk`9I$XSs^T*SU{PfUmznS;>57M>3Pc5lxvmu;_(b8VLdRhV@D_I0?N zdYTB27EcN&n+4$I_5z-7(i3N8Xf_-C6mR=50Zo!A9aPg2W{-ojD|9%gF23bZ@+`c6!6fCG*~;0(I>`U~MRE3F(O6TaR*ofC(X z8#mh$E-Ys0_I+?HA7a2G14`P7mhcIL<>)*Ph?A1$qs?kmJ%)p6B^SM>lr02pX_5Ew z9EJ!PJtnn9NY}EN?d%pp8&VvHs@yf-uobaqMmppGba;rY5z5kD;7fk$_q&)4&Vr+Q z#&|y{6Z3og()v`MEyLn7$R379NW~fY`onal>*S(xM9S#uWV64?JWVk>iUODKu?_$%ss7J>uY|t7sp~eSc?|o1pg&a#{#W zkm^N>hCV(Lx_A79)RG>`tgDu7IAsY>$$oxgBV`;8TLOvMiHXwXC{C>*ua+y$gow znY`!hrv?R&oAnd^uO(~Cn1 zoBCt9PDEwsM*B}xBu{2qSRTf_{&pGu;OAAcxZdfIt`h@8f-OYAk$uI;M%al>vV)$1 z&wVHCLAaubnK-Iz^7Ao?-BW`z1 z#aWsvZ94N_TQ5aYa0KOGjD~mKzvsdP;86UiIub|3pw@kV-o0Z!$NE8*FeYvi1NSgT z&xd}$4EeZVXq9~Hg@ZJ@H?CeP*WC0MXOHu$mi;!)L8^bx)&nP@lsFouK91Dw^BUO~ z9RX-2lj)P)7wl>EI}H!l0_12g5l84s@O%_B9JxvFM)-bQa?6AI8YUVbXu)kKVihA^ zc}O$;qA8qimMGkN1-f(|H@DWb(Pyo39%q@4{4|Ylc!>+>5bfeBA)}DnPC~^Qqs9>J zj_=p-mIS`#ZkRIERX}a@8^{|ALbHH~D$d@kWD~Hwueqr|Pd!@<{kTQCyCxXd-F3$O zU|I#z#<-{n2gTTBY5HFBuy~oI$|Z`69@oy(6HHqzW6>yk}2>W zIzM3Nfb&BGb6pK8K29y$ABZPv}fg%hGFJIVnCuI4||B=Rut7OZu3EYU- z5(NXE29klF@0Fp~K3WMgBNJ<%k8h}G*?D|>UlTXX0bqJl4w^PNmu6ms!RfHQskeSM zudH$??)mcff^Bfzb{ZV=A%{Ws7!!8EZU+Ubk#gV2TikFK9-(elc{}o(kzJFjB$|sc zEQ*8jfnFInou^|?NE5Z=#cGH%2ku2leSkOTD&>;bRB#T_8zqPeey=q=ALPYpd%44c zaveui>>ge(8c5ltFOT)=&f&UF{TP*8k)x@FR-H*<=u71r?UxMNQF9qXMMwoFj!ho? zDP1U|PiS(G?j||I-W65l-V-uZlsKgdyStm@DQZhK_at;@;N+<|D_)fgKY8JV7b2fd!bVwqK^Sp+KQUQz`kao=D`VNI}waNJyJ2qvdbOhx@^f#pe9;9GFkT$Akt=4gtXBZt`IVEv|Mph9iZ?6v5o)!S><_) z$4*6#6}1(gRG%d(G|qr*Kv2P@JG?de9+yh^!RvQf^Mb~|o9?LZ9WwHGqe+W}H3?6V zV|cs`&|8W+sHka%;~MSTZwbBlJ=`Or+LMUP-=shP#dZde$Dcocy1GoJ?0(OaeDf9f zR#+MOBvIKWaK?!hXP(zh?pasTQ?fT7x4^->}0Wc8zaxA@=saaC;(Q1_MWKgx15IDK`h)ep)uY;AY9f0V>R9rN7A zOSO92XrVk@RK_5g-s%bY62c-e=?#rd2097Y@rwy@-92YN*7@Dxr+RUb#r5^tNB#Z% zTYV9c5g81Nq#z9lgqi-%I_JN9f)LJlm-Sclg5R>saWE-ub+QgF+7O53@hH5v69` zB&B89D;ajahKJgQAU6(!K;HY+$;J&f|6Ko?%tH7&#(lGH<1LV4&+tge{bY_Ec7v-0 zlkmb(mB!QHgd7twnHMi!n3&MRR+4B&#Ayb(l0c28DODM|;wx-GmsV40skKsNpQ`va zcPwYBHW>>(83*S=-@ns5SyDy>%*6d5(0WZdo06zfP>QGqBMh8l+qiL?)92y!bNQ-n zP{;D)CI0Enlm*|}F;*k}_WG&b{{3N*WA)oqYkF2(2&i<41Vjy<(K_1V`Af{-4U2d% zgX?!cu3iI;{oEo*^@`K6((hdF)cOfKeZ~D17F(&G&6om%WSRswKnOVE z?o%8nm^1*^5q>*#o6WZ@GqtvG3~JU#SGu(s7Wp7sfEQnW%5>YiM_j-`RacZ>4dR*W zuWEx+~go-3<+{p?!294qUBrlA)2gbI5O)}@cbpSJMt9%CqFc1gap-gv?njc z^yG16o>@L5`nTj9RR4`dI$pAU%_j9o*)H)$GafX`oE+y_-_!Qu3<-<%FYf1UH&sW3 z=i?nj*v;^L+$LXNQVT8jnbqEmByZ}|)$}=A9@nY4Uy?uNXfC0pVo!q_QiiewScvBw zKN_6}--c7gyc)oD6`*)N#f0Gz7I5Mq4USm6JSP;nIe{;`Y`&!DQ#u^r7cqO_I;B1; zZ6e=lhZt<>Vba@qWLy~#o2pAPZ?;u^vR=|}x9VFg%I-RUCBRD`pW*IaKY4YFdvN784$)554ywlvJ6>QepT$#blLDgA7nkSbe0RT< z*YMrng!`;blMCRAhnLklx=1H-6}*3Mak)>yh-S3UtX$V*XdGLKG&)5O(;KAH`Jx1= zV{kY)mxjk=;<{LB(iJ>Olp~FgI6l!Q+!)F4?Ko5AB=%TK(O}o3(i8_E;?1%9=;A=V zyFkez1yr5?DT?J+Kx~e(H>WKB6dG-;b>Ad|QjbP|aK6b~`ZZenyn#>Uj`~8RM$&~AxF8%j){725dg>!bk#M~p_Sp_Q3a1D*ArD!2?b>{j_ z@hl0K?Jbwtq9D~shR(sXO?kGJSh(9_o5z0LeKXuo{f)|B5>SyT*12Q#5n(V5RbYB8 z1qG&wbc`!kHK~nh>Fd~VZh5eKEDi4Rn6)mAJ}bUCd1&;6Z!^Y<*exXDUFBPOA(8fc zVed(Aj6GtPgu>9KzUX6P+uF}E)jEMlwL0-4bad}ArRfOba+a*)=f|KjMsCZ^WY)l_ z)W)8CNT+r?`f}P|Y{yx3dPY=7@qc$@c=0OE+*Ss#w|91Yinp}C!vibMsBe_uE>i2? zszR&5ZHQ5so+yR{NZ`(#T~-jjV-)ZHx#Zbm^yOV|_NMoZIaD$HSy}z>*oAuDh}=nl zeX-{ogg$KVYPve=-x$hi@j99I{XRPFCw8^u|E*AYyykzqP)5Hvu-iT1nz4){bx}KX z+wH;Pasv?#0-0zWWN<)ol@nSl* zw%z=4d-sLkY^X^ddxOI>%cL~))onKuqmTz*zia8j)!RbVAre^lGNJ^NtRsz-D!1a| z(R!w}nrFpo1{%HRB+E12Tw~54g*dXI4`=F|n9Rn#f1$JO-?JUzR_w*d;>p%e(yEss3IcwUg=UT+@=xUA(O(?xJ5m1{$Jgg5-5x*gjqAs|?i6 zZ|`n#oH@VDwm?F<_Y3h|fU14iA=_zlmi<86{lj1Z@GqIKJ zDMjp1l4qDiym#5TZ83b~gk}z8#`51Pc4=^|o?)Z2r#AHl-lY$mll#DDKbhW`m$Hw?@*kTSi;V0Tu#`s zhd*HWO>xcVXybLf^BG0dthCQj3{}c=jQJ#@gg{{2j5XCYBL?v?09|RcZdhIzl}-Y6 z^J(;vwg1-(iWWVAi^q1KM`gSr{H75{uN8Bi3?AX?G#@n(evba--FMd;lin9^nO0iF z>>`^!YKTbI%XlO%q&!s$HK?*{wmoltL@x7CP266i@YcF;-`dFDqOiP{=T^zjz+Hdg zhn&Bv=nWnAt-X;I9L}0$@LI#=iSF}zQTIaVL-X2~9sHJH^^4gouFK|6zSt@CuW^{* zjLg}G{(g0S#92e`UhC1`N_&B+9C8cwmrM+7Ng8j-?{cKL+Zj0fv&#B#q)0K1O=q1~ zL#D_11^rQZ(?=0J3n-dAL>}Wq82O7Dhgl)(%eWiz8h)!xRzsQu@7t<2Qa&cdqj|Ri zt?EeVE9L|ZkVgtc=3r}t9?;&4iCmJ$TL{|K``q>uV^u#5e1?I+X7%*+v@m+@bMVZG zWbfX5`r9SxP!x{bRbm*V$v0F)M0!b0%ugc>_2v}2P!it7N^m#)Qu2R#O7mi=%dq4Q`>#V$L@?7&>8sSHiY+B0m z_3Y@Ye@hK2u^O9gNZNk4OKOMTE?{;!m;Y#Y==ypXck9VXt#!UpHM6xGJL-D@%4dig zoJ#va00RnS5={OoVUNC}e>vyf$CUjI*FZ?5{4|A9guCQ7hrd~6JxO%aX;ZxNn`58c zFrR9>d1b!aY5rA0``O1!k9BgR9gBve)xER7B}-cLY)Mx%M&1g$jWPG=YXk@h&?lJa z=~r70_oc8q#xGfUi{XF0-lj-z@B67bch_ksVpR6MK6wMfOYz&)I}W`o z^JnhNwfYZoDh$l)c4Mvh29?&{amA~MQi;2*srOQKSrtqPH1Aog$W={@%^Tb{8qOrI z`C14`tYy3V-^7cq-(Idl_*YBNzPQ8^7^gc9uT*d#y{6PHn zY2{4bJkGc8fAH70i*NAM4}e&I{F6?&_S)K-*R;+K*G5--**!#&?(oUXJm~csNi?PeA1L7zp{CaBr|-DW z`FPr6bhVq3n`R)LOVWH+CqHJX*Zon9y^T#WA!!CECkEeO8CK`39{>mWX{=g7+UdS* zVNq`jgE}--+IE#QX4qXx6R9JY?-&yaxNbk!#5_qX&)78Z z)<$Li5@Q1-aISsSN2=_$?NEwEw#RxE##xh4qW3$^BpH+v1fMUHu4-m0{}#de+HFlv zvF;6$AMc1ba$%jslt-EZ#EnPj2S?6{lwwegN;Hy*|~|GWu)q4di0V7uDe-v zjBz!Fez@}T4A1*a#=p5SU6Y!(pU={4Xs2b(y$mCdVG1ja*X_ZF?3b(1x zJZmOq4BPdvsg3c6AmN0q(e7#zGYPl?=~N}2A!{K>)q#m6D&2rjZ9g^i5v^+NObkm zWYkpN75A%*m>U3VWJt&YUlsUDRtYj)pnO3d+H~4B%dEa`hAmqB#RI_<+t}1S7)cTjc@qWw!dv0kD8T4sq)` z3tthulK}YGhrdU66%e3L1@F)xGb13Q@LRL46p~R4!*o+;RP-060#O-iV;l_Af1@W4 z87N+@YN5J5AqsMroiLel4YO25X9+GI4rlBHIc(HO=cCA*KKgnlUUepSc_c8DFt)A& zteUhw5}N}Wl!B*LIGQk=PA?U))YsRS9-b}_sbg&AA*nqMJk)g}-xQ{E41^CBC0 z6;oKsRrTaK&&#%V?%zoJSDp+-T$bTiE7L~SFudtXL;fA)yZs&335fX*ZvK$;F^|Ob zZvOFcL*aR@Fo-pk#8H23^r|KDYy9MdMH$Eh`mr-O$tDOeWYerzd@pugOdg@9Odkl* zskM_{{%LSq&%p2teE&up%W5o%^7U|aWfbxM>*c(nn%cH7j1)nNARGvxgQ9=|q+{?P zC{-zfh_nb2S_Dy%P?SScs#IyA2$5!_#n1($Cq|lpQ4VmVgf2Y-7cjzI;lADXd*8Cg zSbL0Qk3BQzUjO(1BrDXyugSyKx5`U~3Hnf*5r>=`PtDLnBAW6${SQKKil}(p;{v=F z$HwnadUSyDvEcixLY+Z~SQsL)-jxlGByqz+Q{yr;?i3fUTX#4vQN9x%OzOg{+#yk$ zk-9Th$Y*vONJAJ8RYSIV%JA6Cn2hhd)c7hG=?y`(-^;RN0{7b4-m zjkot5;&a8GiHx3Lm42xxg-`8=eJ>1z!Q^!@8mp(hRonzrHc9cXM>(DBo)~66}IdM$a)S zw;NP$(*OUb82;dz9R0Dp4CcSBfr(DQgPP#UKq-Ff5bg?2(}%u`odu%+uQyPf&uX^2 zn3vy@WAP?*`1JeVXEqWch#JR4xLm(?(S!46MlI$KCZ%j&80^9r>{AV*J24tpEO8Q8dJ11@dO)Kf{8R#M0ZFZ!)v8 z@Q7xiX6?VE1^mda-YCej&g7FicVisB0lUu<+(oi`!2Zp5wAvj&KbvL1D6T#6+Gim! zARr%(E^BKGTlt#+%r>@7EURGZF#YuHo9!9UIABHIC{(siU{!`uMV9^y{1+G@b0*&1 zz3P8jMXl(C>B&}(^0KmOKbljp8s%X5MwlgzdK$cCQ=u+>SwBa;thXAH2rU`vx-h_J zmjhzK`ft2DuGiFK=Y5t%*s`(W^ux7$jBRnS+(7ILhkl{nlWFRz93%wGKG zTL3^BHd0j{{ICFlEZ^>0i74%SU7Z|Ow~?&tOBjW!d4PDcyOo~l9wE%Rxw%7;As@-? z)z#I0Go(b?DPp_yV?Zt#Lqpo#Fo2n;-b4m$Ee-*4=Xn>)vq*Jt@cEe99cO!-ZADcq z*2CPKXC>Xige(SB2(~S*nMz( zeCPWej^ZnEvq*zy+JHNYI;k$HwLo!e8anGVkFNdr!VImMukDXxPNWxU(oLM$h>U!PeMLcps`f>2za-ULQeVoQ`%`fM(bRO1 z^*}ERnc2O^sU>!Dy25<*v0%;n8;i7`lcowikIG2owe^)o`akd@r7n|-UaXUY2nw+n zchDnUv9#eW-jq<17t+3AW#GW&%@rqrZ=qegR+!hQAqFVb9uRzIkOk(j=Xj{wl_eL2 zm5&dx6ji=dZhLLy)Q>HUhIi(U>mTm8xO7lFHGp0(6ikc*>@)HmyqSw0&~Tycm5QBH_7BHA^Xq2(YoPc;KhA$Xy&Vh)IU=4lb$)y4FfECZ)Ux1 z8fZWj&2K(erAfVn80C8LoRfL}i2_qy{k>|hK9iOewrvn*+3>bVJ@BWD*Pso=`_7Q+ zr;#WL+xw)Kg(4{mh^JJqt)^i+Sr3ubm5C3{4GmiJL+Yy2YvHWd4NX{D&tZu&uXeet zGV5xHnVV5%%U&kQINmKAHN|=XJZ9CCT$e;TxV5;AkKVW573BZi>Ag2iA51K;v?kH+ znx)XC+RF?>1MI4Pj7LU*QwE=z%KPq}j^kDIxwe%izimXeoO}ZPV^Ns@44K;P$*N1x z{)D|r8zS<;=?{~tWL~}|zSAb2yWC{gB35xfvDfh?p*5)io=>kC{CYk|QV9mFpItH` z2KAffglFn!B~AkWB)8KggQQ1yV_j2xU@6Qm;&6(M>bz2`<^{1mTFN@Vy-v}EdFB^7I89O{R3*tGw{r|SXMSD>X{eYj zr4mXHS|!`Tz6K$K)ogPtp4bU&QiaP!50Vz-3dvsalX zF46mjZKvrSrfdf}0<*wybY1-yQX&3ezawst`m(SiEGjkTa_Dtyt4tw+LpNUSeHghh zUVS@8uvVxGr2zL;NSpATh!yQ!JLIgjJIol^v%Xdc`&~wHg*c1uDk@J1r`@_BQE)ys zXSVA@EFKEgZOt@``Z58qTb4e0pk^%XRp125Dt$DC13W-qj;uqx&(b zK_?#tp`_&ad}sJ-Wqw?#ON?TasHa(Px=MG+Q zw)dj8{jJpFTAY97nnd^PCH-S&oEB&DkfoAt*(V3Rx4Rk(S2c29)kDhXfqpW$oG{*_ z7+?M6ENPv$e146^@sZfSK9xxKY>E!dovO&O%+o%`^{^lhZ8cCr%wH5OpRf6R(k}kj z?8+Var_A9`6N;6ZJ(awUGJClBu#(Rfj`EcL#~Gfr;9o49F@&=-hNN}`__&tC>i}e1 ky`xkt#<&8Q#Y4%l_p-OA{#qyl>8MPmh86}@S6m|h4Kg$#EC2ui literal 0 HcmV?d00001