From 2923721108b365625a35779956f3f0277f446401 Mon Sep 17 00:00:00 2001 From: lyifa Date: Tue, 10 Sep 2024 22:21:26 -0400 Subject: [PATCH 01/14] CPU Compaction --- stream_compaction/cpu.cu | 34 ++++++++++++++++++++++++++++++---- 1 file changed, 30 insertions(+), 4 deletions(-) diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 719fa11..635c98f 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -18,9 +18,13 @@ namespace StreamCompaction { * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); + //timer().startCpuTimer(); // TODO - timer().endCpuTimer(); + odata[0] = 0; + for (int i = 1; i < n; i++) { + odata[i] = odata[i - 1] + idata[i]; + } + //timer().endCpuTimer(); } /** @@ -31,8 +35,15 @@ namespace StreamCompaction { int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int count = 0; + for (int i = 0; i < n; i++) { + if (idata[i] != 0) { + odata[count] = idata[i]; + count++; + } + } timer().endCpuTimer(); - return -1; + return count; } /** @@ -43,8 +54,23 @@ namespace StreamCompaction { int compactWithScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); // TODO + int *bools = new int[n]; + for (int i = 0; i < n; i++) { + bools[i] = (idata[i] != 0) ? 1 : 0; + } + int* scanResult = new int[n]; + scan(n, scanResult, bools); + // scatter + int count = bools[n - 1] == 1 ? scanResult[n - 1] : scanResult[n - 1] + 1; + for (int i = 0; i < n; i++) { + if (bools[i] == 1) { + odata[scanResult[i]] = idata[i]; + } + } timer().endCpuTimer(); - return -1; + delete[] bools; + delete[] scanResult; + return count; } } } From ef7b7483452f4862e4a2fc569e826dea43954194 Mon Sep 17 00:00:00 2001 From: lyifa Date: Sat, 14 Sep 2024 11:36:07 -0400 Subject: [PATCH 02/14] Navie GPU scan and fix CPU scan bug --- src/main.cpp | 4 ++-- stream_compaction/cpu.cu | 4 ++-- stream_compaction/naive.cu | 41 ++++++++++++++++++++++++++++++++++++++ 3 files changed, 45 insertions(+), 4 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 896ac2b..1c5ac84 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -51,7 +51,7 @@ int main(int argc, char* argv[]) { printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(SIZE, b, c); /* For bug-finding only: Array of 1s to help find bugs in stream compaction or scan @@ -64,7 +64,7 @@ int main(int argc, char* argv[]) { printDesc("naive scan, non-power-of-two"); StreamCompaction::Naive::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 635c98f..fd31c1d 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -22,7 +22,7 @@ namespace StreamCompaction { // TODO odata[0] = 0; for (int i = 1; i < n; i++) { - odata[i] = odata[i - 1] + idata[i]; + odata[i] = odata[i - 1] + idata[i - 1]; } //timer().endCpuTimer(); } @@ -61,7 +61,7 @@ namespace StreamCompaction { int* scanResult = new int[n]; scan(n, scanResult, bools); // scatter - int count = bools[n - 1] == 1 ? scanResult[n - 1] : scanResult[n - 1] + 1; + int count = bools[n - 1] == 1 ? scanResult[n - 1] : scanResult[n - 1]; for (int i = 0; i < n; i++) { if (bools[i] == 1) { odata[scanResult[i]] = idata[i]; diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 4308876..0351282 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -12,6 +12,32 @@ namespace StreamCompaction { return timer; } // TODO: __global__ + int getLog2(int n) { + int log2 = 0; + while (n >>= 1) { + log2++; + } + return log2+1; + } + + __global__ void kernScan(int n, int* odata, const int* idata, int log2_n) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + // exclusive scan + odata[index] = (index > 0) ? idata[index - 1] : 0; + __syncthreads(); + + for (int d = 1; d <= log2_n; ++d) { + int t = 1 << (d - 1); + int temp = 0; + if (index >= t) temp = odata[index - t]; // Load the previous step's value + __syncthreads(); // Synchronize before updating + if (index >= t) odata[index] += temp; // Update the current value + __syncthreads(); // Synchronize after updating + } + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. @@ -19,6 +45,21 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata) { timer().startGpuTimer(); // TODO + int log2_n = getLog2(n); + //printf("n: %d\n", n); + //printf("log2_n: %d\n", log2_n); + // call kernel + int* dev_idata; + int* dev_odata; + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + kernScan << <1, n >> > (n, dev_odata, dev_idata, log2_n); + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaFree(dev_idata); + cudaFree(dev_odata); + + timer().endGpuTimer(); } } From 251e453983aad763490000f3504bd6dec7d9332b Mon Sep 17 00:00:00 2001 From: lyifa Date: Sun, 15 Sep 2024 12:08:47 -0400 Subject: [PATCH 03/14] Work-Efficient GPU Scan --- src/main.cpp | 6 +- src/testing_helpers.hpp | 4 +- stream_compaction/efficient.cu | 110 +++++++++++++++++++++++++++++++++ 3 files changed, 115 insertions(+), 5 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 1c5ac84..c47cba3 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 << 3; // 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]; @@ -71,14 +71,14 @@ int main(int argc, char* argv[]) { printDesc("work-efficient scan, power-of-two"); StreamCompaction::Efficient::scan(SIZE, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(SIZE, c, true); + printArray(SIZE, c, true); printCmpResult(SIZE, b, c); zeroArray(SIZE, c); printDesc("work-efficient scan, non-power-of-two"); StreamCompaction::Efficient::scan(NPOT, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(NPOT, c, true); + printArray(NPOT, c, true); printCmpResult(NPOT, b, c); zeroArray(SIZE, c); diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index 025e94a..af9a1ff 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -61,8 +61,8 @@ void printArray(int n, int *a, bool abridged = false) { printf(" [ "); for (int i = 0; i < n; i++) { if (abridged && i + 2 == 15 && n > 16) { - i = n - 2; - printf("... "); + //i = n - 2; + //printf("... "); } printf("%3d ", a[i]); } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 2db346e..f69e479 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,12 +12,122 @@ namespace StreamCompaction { return timer; } + // up-sweep kernel + __global__ void kernUpSweep(int n, int* odata, const int* idata, int t) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + // exclusive scan + odata[index] = (index > 0) ? idata[index - 1] : 0; + __syncthreads(); + // upsweep + for (int d = 0; d <= t; ++d) { + int offset = 1 << (d + 1); + int ai = index + offset - 1; + int bi = index + (offset / 2) - 1; + if (index < n && (index % offset) == 0) { + odata[ai] += odata[bi]; + } + + __syncthreads(); + } + } + + // down-sweep kernel + __global__ void kernDownSweep(int n, int* odata, const int* idata, int t) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= 1 << (t + 1)) { + return; + } + // exclusive scan + odata[index] = (index > 0) ? idata[index - 1] : 0; + __syncthreads(); + // downsweep + if (index == 0) { + odata[n - 1] = 0; + } + for (int d = t; d >= 0; --d) { + int offset = 1 << (d + 1); + int ai = index + offset - 1; + int bi = index + (offset / 2) - 1; + if (index < n && (index % offset) == 0) { + int temp = odata[bi]; + odata[bi] = odata[ai]; + odata[ai] += temp; + } + + __syncthreads(); + } + } + + // up sweep + down aweep + __global__ void kernScan(int n, int* odata, const int* idata, int t) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + int paddedSize = 1 << (t + 1); + if (index >= paddedSize) { + return; + } + // exclusive scan + //odata[index] = (index > 0) ? idata[index - 1] : 0; + //odata[index] = idata[index]; + odata[index] = (index >= n) ? 0 : idata[index]; + __syncthreads(); + // upsweep + for (int d = 0; d <= t; ++d) { + int offset = 1 << (d + 1); + int ai = index + offset - 1; + int bi = index + (offset / 2) - 1; + if (index < paddedSize && (index % offset) == 0) { + odata[ai] += odata[bi]; + } + + __syncthreads(); + } + // downsweep + if (index == 0) { + odata[paddedSize - 1] = 0; + } + + for (int d = t; d >= 0; --d) { + int offset = 1 << (d + 1); + int ai = index + offset - 1; + int bi = index + (offset / 2) - 1; + if (index < paddedSize && (index % offset) == 0) { + int temp = odata[bi]; + odata[bi] = odata[ai]; + odata[ai] += temp; + } + + __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 + // call kernel + int* dev_idata; + int* dev_odata; + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + int t = ilog2ceil(n) - 1; + printf("log2_n - 1: %d\n", t); + //kernUpSweep << <1, n >> > (n, dev_odata, dev_idata, t); + //kernDownSweep << <1, n >> > (n, dev_odata, dev_idata, t); + kernScan << <1, n >> > (n, dev_odata, dev_idata, t); + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaFree(dev_idata); + cudaFree(dev_odata); + + + + timer().endGpuTimer(); } From f0dcd5bbd8ca2b50f93e3795890a1b942f43ac6b Mon Sep 17 00:00:00 2001 From: lyifa Date: Sun, 15 Sep 2024 12:48:59 -0400 Subject: [PATCH 04/14] Work-Efficient Stream Compaction --- src/main.cpp | 4 ++-- stream_compaction/common.cu | 14 ++++++++++++++ stream_compaction/efficient.cu | 35 ++++++++++++++++++++++++++++++---- 3 files changed, 47 insertions(+), 6 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index c47cba3..27b9997 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -137,14 +137,14 @@ int main(int argc, char* argv[]) { printDesc("work-efficient compact, power-of-two"); count = StreamCompaction::Efficient::compact(SIZE, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(count, c, true); + printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); zeroArray(SIZE, c); printDesc("work-efficient compact, non-power-of-two"); count = StreamCompaction::Efficient::compact(NPOT, c, a); printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)"); - //printArray(count, c, true); + printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); system("pause"); // stop Win32 console from closing on exit diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu index 2ed6d63..bd5071c 100644 --- a/stream_compaction/common.cu +++ b/stream_compaction/common.cu @@ -24,6 +24,12 @@ namespace StreamCompaction { */ __global__ void kernMapToBoolean(int n, int *bools, const int *idata) { // TODO + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + bools[index] = (idata[index] != 0) ? 1 : 0; + } /** @@ -33,6 +39,14 @@ namespace StreamCompaction { __global__ void kernScatter(int n, int *odata, const int *idata, const int *bools, const int *indices) { // TODO + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + if (bools[index] == 1) { + odata[indices[index]] = idata[index]; + } + } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index f69e479..911143b 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -125,9 +125,6 @@ namespace StreamCompaction { cudaFree(dev_idata); cudaFree(dev_odata); - - - timer().endGpuTimer(); } @@ -143,8 +140,38 @@ namespace StreamCompaction { int compact(int n, int *odata, const int *idata) { timer().startGpuTimer(); // TODO + // compute bool array + int* dev_bools; + int* dev_idata; + int* dev_indices; + int* dev_odata; + int* bools = new int[n]; + int* indices = new int[n]; + cudaMalloc((void**)&dev_bools, n * sizeof(int)); + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + cudaMalloc((void**)&dev_indices, n * sizeof(int)); + cudaMalloc((void**)&dev_odata, n * sizeof(int)); + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + StreamCompaction::Common::kernMapToBoolean << <1, n >> > (n, dev_bools, dev_idata); + // scan + int t = ilog2ceil(n) - 1; + kernScan << <1, n >> > (n, dev_indices, dev_bools, t); + // scatter + StreamCompaction::Common::kernScatter << <1, n >> > (n, dev_odata, dev_idata, dev_bools, dev_indices); + cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(bools, dev_bools, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(indices, dev_indices, n * sizeof(int), cudaMemcpyDeviceToHost); + int count = bools[n - 1] ? indices[n - 1] + 1 : indices[n - 1]; + cudaFree(dev_bools); + cudaFree(dev_idata); + cudaFree(dev_indices); + cudaFree(dev_odata); + delete[] bools; + delete[] indices; + + timer().endGpuTimer(); - return -1; + return count; } } } From 80ae481335b0826b54def06e4a15b2172838738c Mon Sep 17 00:00:00 2001 From: lyifa Date: Sun, 15 Sep 2024 13:23:36 -0400 Subject: [PATCH 05/14] thrust scan --- src/main.cpp | 2 +- src/testing_helpers.hpp | 4 ++-- stream_compaction/thrust.cu | 5 +++++ 3 files changed, 8 insertions(+), 3 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 27b9997..98f2ccf 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 3; // feel free to change the size of array +const int SIZE = 1 << 30; // 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/src/testing_helpers.hpp b/src/testing_helpers.hpp index af9a1ff..025e94a 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -61,8 +61,8 @@ void printArray(int n, int *a, bool abridged = false) { printf(" [ "); for (int i = 0; i < n; i++) { if (abridged && i + 2 == 15 && n > 16) { - //i = n - 2; - //printf("... "); + i = n - 2; + printf("... "); } printf("%3d ", a[i]); } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1def45e..8546b5d 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -22,6 +22,11 @@ namespace StreamCompaction { // 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::device_vector dv_in(idata, idata + n); + thrust::device_vector dv_out(n); + thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + thrust::copy(dv_out.begin(), dv_out.end(), odata); + timer().endGpuTimer(); } } From 21832a82b19aac9e0e87b766e34fbb8ce70c7907 Mon Sep 17 00:00:00 2001 From: lyifa Date: Mon, 16 Sep 2024 00:07:13 -0400 Subject: [PATCH 06/14] shared memory and arbitrary length for naive sollution --- src/main.cpp | 2 +- stream_compaction/efficient.cu | 10 ++- stream_compaction/naive.cu | 110 ++++++++++++++++++++++++++++++--- 3 files changed, 109 insertions(+), 13 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 98f2ccf..6afe591 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 30; // feel free to change the size of array +const int SIZE = 1 << 4; // 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 911143b..c7993a5 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -108,8 +108,10 @@ namespace StreamCompaction { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); + // TODO + const int blockSize = 128; + dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); // call kernel int* dev_idata; int* dev_odata; @@ -118,14 +120,16 @@ namespace StreamCompaction { cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); int t = ilog2ceil(n) - 1; printf("log2_n - 1: %d\n", t); + timer().startGpuTimer(); //kernUpSweep << <1, n >> > (n, dev_odata, dev_idata, t); //kernDownSweep << <1, n >> > (n, dev_odata, dev_idata, t); - kernScan << <1, n >> > (n, dev_odata, dev_idata, t); + kernScan << > > (n, dev_odata, dev_idata, t); + timer().endGpuTimer(); cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_idata); cudaFree(dev_odata); - timer().endGpuTimer(); + } /** diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 0351282..443350c 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -21,13 +21,14 @@ namespace StreamCompaction { } __global__ void kernScan(int n, int* odata, const int* idata, int log2_n) { - int index = threadIdx.x + (blockIdx.x * blockDim.x); - if (index >= n) { + int index = (blockIdx.x * blockDim.x) + threadIdx.x; + int pedding = 1 << log2_n; + if (index >= pedding) { return; } // exclusive scan odata[index] = (index > 0) ? idata[index - 1] : 0; - __syncthreads(); + __syncthreads(); // odata with first element as 0 is ready for (int d = 1; d <= log2_n; ++d) { int t = 1 << (d - 1); @@ -39,28 +40,119 @@ namespace StreamCompaction { } } + __global__ void kernBlockWiseExclusiveScan(int n, int* odata, const int* idata, int blockSize) { + extern __shared__ int sdata[]; + + int idx = threadIdx.x; + int blockStartIndex = blockIdx.x * blockDim.x; + int index = blockStartIndex + idx; + + // Load data into shared memory + if (index < n) { + sdata[idx] = idata[index];//(idx > 0) ? idata[index - 1] : 0; + } + else { + sdata[idx] = 0; // Out-of-range threads + } + __syncthreads(); + + // Perform in-block scan + for (int d = 1; d < blockDim.x; d *= 2) { + int t = idx >= d ? sdata[idx - d] : 0; + __syncthreads(); + if (idx >= d) { + sdata[idx] += t; + } + __syncthreads(); + } + + // Write results to global memory + if (index < n) { + odata[index] = sdata[idx]; + } + } + + + + // kernel for write total sum of each block into a new array + __global__ void kernWriteBlockSum(int n, const int* odata, int* blockSum) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + + if ((index + 1) % (blockDim.x) == 0) { + int i = (index + 1) / (blockDim.x) - 1; + blockSum[i] = odata[index]; + } + } + + // kernel for add block increments to each element in the corresponding block + __global__ void kernAddBlockSum(int n, int* odata, const int* blockSum) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + if (index >= n) { + return; + } + + // exclusive scan + // set first element to 0 + if (index == 0) { + odata[index] = 0; + } + else { + // Add the block sum from the previous blocks to the current element, except the very first element + odata[index] = odata[index - 1] + blockSum[blockIdx.x]; + } + } + + + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO + // TODO int log2_n = getLog2(n); + int t = ilog2ceil(n); //printf("n: %d\n", n); - //printf("log2_n: %d\n", log2_n); + printf("log2_n: %d\n", t); + int blockSize = 4; + int numBlocks = (n + blockSize - 1) / blockSize; + dim3 fullBlocksPerGrid(numBlocks); + printf("block size: %d\n", blockSize); + printf("numBlocks: %d\n", numBlocks); // call kernel int* dev_idata; int* dev_odata; + int* dev_blockSum; + int* dev_blockIncrements; + int* blockSum = new int[n]; cudaMalloc((void**)&dev_idata, n * sizeof(int)); cudaMalloc((void**)&dev_odata, n * sizeof(int)); + cudaMalloc((void**)&dev_blockSum, n * sizeof(int)); + cudaMalloc((void**)&dev_blockIncrements, n * sizeof(int)); cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); - kernScan << <1, n >> > (n, dev_odata, dev_idata, log2_n); + cudaMemcpy(dev_blockSum, blockSum, blockSize * sizeof(int), cudaMemcpyHostToDevice); + timer().startGpuTimer(); + // scan on each block + //kernScan << > > (n, dev_odata, dev_idata, t); + kernBlockWiseExclusiveScan << > > (n, dev_odata, dev_idata, blockSize); + // write total sum of each block to blockSum + kernWriteBlockSum << > > (n, dev_odata, dev_blockSum); + // scan on blockSum + int blockSumSize = ilog2ceil(numBlocks); + kernScan << <1, numBlocks >> > (numBlocks, dev_blockIncrements, dev_blockSum, blockSumSize); + // add block increments to each element in the corresponding block + kernAddBlockSum << > > (n, dev_odata, dev_blockIncrements); + //dev_odata = dev_blockIncrements; + timer().endGpuTimer(); cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_idata); cudaFree(dev_odata); + cudaFree(dev_blockSum); + delete[] blockSum; - - timer().endGpuTimer(); + } } } From d6a8fee19c73d002f3b337cff11091ad0ed921cb Mon Sep 17 00:00:00 2001 From: lyifa Date: Mon, 16 Sep 2024 00:34:04 -0400 Subject: [PATCH 07/14] fix index bug for arbitrary length --- src/main.cpp | 2 +- stream_compaction/naive.cu | 12 ++++++++---- 2 files changed, 9 insertions(+), 5 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 6afe591..7a15a41 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 4; // feel free to change the size of array +const int SIZE = 1 << 11; // 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/naive.cu b/stream_compaction/naive.cu index 443350c..9ba7ec0 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -96,13 +96,17 @@ namespace StreamCompaction { // exclusive scan // set first element to 0 + + int temp = 0; if (index == 0) { - odata[index] = 0; + temp = 0; } else { - // Add the block sum from the previous blocks to the current element, except the very first element - odata[index] = odata[index - 1] + blockSum[blockIdx.x]; + int blockIdx = (index - 1) / blockDim.x; + int sumToAdd = blockSum[blockIdx]; + temp = odata[index - 1] + sumToAdd; } + odata[index] = temp; } @@ -116,7 +120,7 @@ namespace StreamCompaction { int t = ilog2ceil(n); //printf("n: %d\n", n); printf("log2_n: %d\n", t); - int blockSize = 4; + int blockSize = 256; int numBlocks = (n + blockSize - 1) / blockSize; dim3 fullBlocksPerGrid(numBlocks); printf("block size: %d\n", blockSize); From 36d37b9fc4ea693f4e5372a1830430f2cb7bc170 Mon Sep 17 00:00:00 2001 From: lyifa Date: Tue, 17 Sep 2024 15:22:24 -0400 Subject: [PATCH 08/14] sanity check --- src/main.cpp | 2 +- src/testing_helpers.hpp | 8 +-- stream_compaction/cpu.cu | 8 +-- stream_compaction/efficient.cu | 92 +++++++++++++++++++++++++++++++--- stream_compaction/naive.cu | 51 +++++++++++++++++-- 5 files changed, 141 insertions(+), 20 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 7a15a41..6afe591 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 11; // feel free to change the size of array +const int SIZE = 1 << 4; // 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/src/testing_helpers.hpp b/src/testing_helpers.hpp index 025e94a..8dc5bdd 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -10,7 +10,9 @@ template int cmpArrays(int n, T *a, T *b) { for (int i = 0; i < n; i++) { if (a[i] != b[i]) { - printf(" a[%d] = %d, b[%d] = %d\n", i, a[i], i, b[i]); + printf(" a[%d] = %d, b[%d] = %d\n ", i, a[i], i, b[i]); + printf(" a[%d] = %d, b[%d] = %d\n ", i - 1, a[i - 1], i - 1, b[i - 1]); + printf(" a[%d] = %d, b[%d] = %d\n ", i +1, a[i + 1], i + 1, b[i + 1]); return 1; } } @@ -61,8 +63,8 @@ void printArray(int n, int *a, bool abridged = false) { printf(" [ "); for (int i = 0; i < n; i++) { if (abridged && i + 2 == 15 && n > 16) { - i = n - 2; - printf("... "); + //i = n - 2; + //printf("... "); } printf("%3d ", a[i]); } diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index fd31c1d..8bf1ab6 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -18,13 +18,13 @@ namespace StreamCompaction { * (Optional) For better understanding before starting moving to GPU, you can simulate your GPU scan in this function first. */ void scan(int n, int *odata, const int *idata) { - //timer().startCpuTimer(); + timer().startCpuTimer(); // TODO odata[0] = 0; for (int i = 1; i < n; i++) { odata[i] = odata[i - 1] + idata[i - 1]; } - //timer().endCpuTimer(); + timer().endCpuTimer(); } /** @@ -52,7 +52,7 @@ namespace StreamCompaction { * @returns the number of elements remaining after compaction. */ int compactWithScan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); + //timer().startCpuTimer(); // TODO int *bools = new int[n]; for (int i = 0; i < n; i++) { @@ -67,7 +67,7 @@ namespace StreamCompaction { odata[scanResult[i]] = idata[i]; } } - timer().endCpuTimer(); + //timer().endCpuTimer(); delete[] bools; delete[] scanResult; return count; diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index c7993a5..a062ce2 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -3,6 +3,9 @@ #include "common.h" #include "efficient.h" +#define blockSize 1 +#define RECURSIVE_SCAN 0 + namespace StreamCompaction { namespace Efficient { using StreamCompaction::Common::PerformanceTimer; @@ -12,12 +15,14 @@ namespace StreamCompaction { return timer; } + // up-sweep kernel __global__ void kernUpSweep(int n, int* odata, const int* idata, int t) { int index = threadIdx.x + (blockIdx.x * blockDim.x); if (index >= n) { return; } +#if RECURSIVE_SCAN // exclusive scan odata[index] = (index > 0) ? idata[index - 1] : 0; __syncthreads(); @@ -32,11 +37,23 @@ namespace StreamCompaction { __syncthreads(); } - } +#else + odata[index] = idata[index]; + __syncthreads(); + int offset = 1 << (t + 1); // 2^(d + 1) + int ai = index + offset - 1; + int bi = index + (offset / 2) - 1; + if (index < n && (index % offset) == 0) { + odata[ai] += idata[bi]; + } + __syncthreads(); +#endif + } // down-sweep kernel __global__ void kernDownSweep(int n, int* odata, const int* idata, int t) { int index = threadIdx.x + (blockIdx.x * blockDim.x); +#if RECURSIVE_SCAN if (index >= 1 << (t + 1)) { return; } @@ -59,6 +76,27 @@ namespace StreamCompaction { __syncthreads(); } +#else + if (index >= n) { + return; + } + + __syncthreads(); + + int offset = 1 << (t + 1); + int ai = index + offset - 1; + int bi = index + (offset / 2) - 1; + if (index % offset == 0) { + int temp = odata[bi]; + odata[bi] = odata[ai]; + odata[ai] += temp; + } + + __syncthreads(); + + +#endif + } // up sweep + down aweep @@ -110,20 +148,58 @@ namespace StreamCompaction { void scan(int n, int *odata, const int *idata) { // TODO - const int blockSize = 128; - dim3 fullBlocksPerGrid((n + blockSize - 1) / blockSize); + int t = ilog2ceil(n) - 1; + int peddedSize = 1 << (t + 1); + //const int blockSize = 128; + int numBlocks = (peddedSize + blockSize - 1) / blockSize; + dim3 fullBlocksPerGrid(numBlocks); + + printf("log2_n - 1: %d\n", t); + printf("array size: %d; pedded size: %d\n", n, peddedSize); + printf("block numbers: %d\n", numBlocks); // call kernel int* dev_idata; int* dev_odata; - cudaMalloc((void**)&dev_idata, n * sizeof(int)); - cudaMalloc((void**)&dev_odata, n * sizeof(int)); + cudaMalloc((void**)&dev_idata, peddedSize * sizeof(int)); + cudaMalloc((void**)&dev_odata, peddedSize * sizeof(int)); + cudaMemset(dev_odata, 0, peddedSize * sizeof(int)); + cudaMemset(dev_idata, 0, peddedSize * sizeof(int)); cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); - int t = ilog2ceil(n) - 1; - printf("log2_n - 1: %d\n", t); + + + timer().startGpuTimer(); + +#if RECURSIVE_SCAN //kernUpSweep << <1, n >> > (n, dev_odata, dev_idata, t); //kernDownSweep << <1, n >> > (n, dev_odata, dev_idata, t); - kernScan << > > (n, dev_odata, dev_idata, t); + //kernScan << > > (n, dev_odata, dev_idata, t); // arbitrary block size + kernScan << <1, n >> > (n, dev_odata, dev_idata, t); +#else + + // up-sweep + for (int d = 0; d <= t; ++d) { + int offset = 1 << (d + 1); + //int numBlocks = (n + offset - 1) / offset; + //dim3 fullBlocksPerGrid(numBlocks); + kernUpSweep << > > (peddedSize, dev_odata, dev_idata, d); + int* temp = dev_idata; + dev_idata = dev_odata; + dev_odata = temp; + } + // down sweep + // set last element to 0 + dev_odata = dev_idata; + cudaMemset(dev_odata + peddedSize - 1, 0, sizeof(int)); + for (int d = t; d >= 0; d--) { + int offset = 1 << (d + 1); + //int numBlocks = (n + offset - 1) / offset; + //dim3 fullBlocksPerGrid(numBlocks); + + kernDownSweep << > > (peddedSize, dev_odata, dev_idata, d); + } + +#endif timer().endGpuTimer(); cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_idata); diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 9ba7ec0..31e7750 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -3,6 +3,8 @@ #include "common.h" #include "naive.h" +# define RECURSIVE_SCAN 0 + namespace StreamCompaction { namespace Naive { using StreamCompaction::Common::PerformanceTimer; @@ -20,6 +22,21 @@ namespace StreamCompaction { return log2+1; } + __global__ void kernInclusiveScan(int n, int* odata, const int* idata, int t) { + int index = threadIdx.x + (blockIdx.x * blockDim.x); + + if (index >= n) { + return; + } + + odata[index] = idata[index]; + __syncthreads(); + + if(index >= t) { + odata[index] = idata[index - t] + odata[index]; + } + } + __global__ void kernScan(int n, int* odata, const int* idata, int log2_n) { int index = (blockIdx.x * blockDim.x) + threadIdx.x; int pedding = 1 << log2_n; @@ -125,7 +142,10 @@ namespace StreamCompaction { dim3 fullBlocksPerGrid(numBlocks); printf("block size: %d\n", blockSize); printf("numBlocks: %d\n", numBlocks); - // call kernel + // get block size and block number for scan block sum + int numBlocks_scan = (numBlocks + blockSize - 1) / blockSize; + dim3 fullBlocksPerGrid_scan(numBlocks_scan); + int* dev_idata; int* dev_odata; int* dev_blockSum; @@ -137,20 +157,43 @@ namespace StreamCompaction { cudaMalloc((void**)&dev_blockIncrements, n * sizeof(int)); cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); cudaMemcpy(dev_blockSum, blockSum, blockSize * sizeof(int), cudaMemcpyHostToDevice); + timer().startGpuTimer(); + // call kernel +# if RECURSIVE_SCAN // scan on each block - //kernScan << > > (n, dev_odata, dev_idata, t); - kernBlockWiseExclusiveScan << > > (n, dev_odata, dev_idata, blockSize); + //kernScan << > > (n, dev_odata, dev_idata, t); // non-shared memory one + kernBlockWiseExclusiveScan << > > (n, dev_odata, dev_idata, blockSize); // shared memory one // write total sum of each block to blockSum kernWriteBlockSum << > > (n, dev_odata, dev_blockSum); // scan on blockSum int blockSumSize = ilog2ceil(numBlocks); kernScan << <1, numBlocks >> > (numBlocks, dev_blockIncrements, dev_blockSum, blockSumSize); + //kernScan << > > (numBlocks, dev_blockIncrements, dev_blockSum, blockSumSize); + // recursive scan on blockSum // add block increments to each element in the corresponding block kernAddBlockSum << > > (n, dev_odata, dev_blockIncrements); - //dev_odata = dev_blockIncrements; + //dev_odata = dev_blockIncrements; // for testing timer().endGpuTimer(); cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); +# else + for (int d = 1; d <= log2_n; d++) { + int pedding = 1 << (d - 1); + kernInclusiveScan << > > (n, dev_odata, dev_idata, pedding); + int* temp = dev_idata; + dev_idata = dev_odata; + dev_odata = temp; + } + + + timer().endGpuTimer(); + + // right shift odata + odata[0] = 0; + cudaMemcpy(odata + 1, dev_idata, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); + +#endif + cudaFree(dev_idata); cudaFree(dev_odata); cudaFree(dev_blockSum); From a4f618c4bd680e8b942176049b7f604820c49b9f Mon Sep 17 00:00:00 2001 From: lyifa Date: Tue, 17 Sep 2024 16:14:47 -0400 Subject: [PATCH 09/14] fix work-efficient up-sweeping bug --- src/main.cpp | 2 +- src/testing_helpers.hpp | 4 +-- stream_compaction/efficient.cu | 60 +++++++++++++++++++++------------- 3 files changed, 40 insertions(+), 26 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 6afe591..8a8fb5a 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -13,7 +13,7 @@ #include #include "testing_helpers.hpp" -const int SIZE = 1 << 4; // feel free to change the size of array +const int SIZE = 1 << 10; // 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/src/testing_helpers.hpp b/src/testing_helpers.hpp index 8dc5bdd..ba0e19f 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -63,8 +63,8 @@ void printArray(int n, int *a, bool abridged = false) { printf(" [ "); for (int i = 0; i < n; i++) { if (abridged && i + 2 == 15 && n > 16) { - //i = n - 2; - //printf("... "); + i = n - 2; + printf("... "); } printf("%3d ", a[i]); } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index a062ce2..ffc52a2 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -3,7 +3,7 @@ #include "common.h" #include "efficient.h" -#define blockSize 1 +#define blockSize 32 #define RECURSIVE_SCAN 0 namespace StreamCompaction { @@ -17,7 +17,7 @@ namespace StreamCompaction { // up-sweep kernel - __global__ void kernUpSweep(int n, int* odata, const int* idata, int t) { + __global__ void kernUpSweep(int n, int* odata, int* idata, int t) { int index = threadIdx.x + (blockIdx.x * blockDim.x); if (index >= n) { return; @@ -38,15 +38,12 @@ namespace StreamCompaction { __syncthreads(); } #else - odata[index] = idata[index]; - __syncthreads(); int offset = 1 << (t + 1); // 2^(d + 1) int ai = index + offset - 1; int bi = index + (offset / 2) - 1; - if (index < n && (index % offset) == 0) { - odata[ai] += idata[bi]; + if (index < n && ((index) % offset) == 0) { + idata[ai] += idata[bi]; } - __syncthreads(); #endif } @@ -165,9 +162,8 @@ namespace StreamCompaction { cudaMemset(dev_odata, 0, peddedSize * sizeof(int)); cudaMemset(dev_idata, 0, peddedSize * sizeof(int)); cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); - - - + cudaMemcpy(dev_odata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + timer().startGpuTimer(); #if RECURSIVE_SCAN @@ -183,25 +179,20 @@ namespace StreamCompaction { //int numBlocks = (n + offset - 1) / offset; //dim3 fullBlocksPerGrid(numBlocks); kernUpSweep << > > (peddedSize, dev_odata, dev_idata, d); - int* temp = dev_idata; - dev_idata = dev_odata; - dev_odata = temp; } // down sweep // set last element to 0 - dev_odata = dev_idata; - cudaMemset(dev_odata + peddedSize - 1, 0, sizeof(int)); + cudaMemset(dev_idata + peddedSize - 1, 0, sizeof(int)); for (int d = t; d >= 0; d--) { int offset = 1 << (d + 1); //int numBlocks = (n + offset - 1) / offset; - //dim3 fullBlocksPerGrid(numBlocks); - - kernDownSweep << > > (peddedSize, dev_odata, dev_idata, d); + //dim3 fullBlocksPerGrid(numBlocks); + kernDownSweep << > > (peddedSize, dev_idata, dev_idata, d); } #endif timer().endGpuTimer(); - cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dev_idata); cudaFree(dev_odata); @@ -218,9 +209,11 @@ namespace StreamCompaction { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - timer().startGpuTimer(); + //timer().startGpuTimer(); // TODO // compute bool array + int t = ilog2ceil(n) - 1; + int peddedSize = 1 << (t + 1); int* dev_bools; int* dev_idata; int* dev_indices; @@ -232,10 +225,30 @@ namespace StreamCompaction { cudaMalloc((void**)&dev_indices, n * sizeof(int)); cudaMalloc((void**)&dev_odata, n * sizeof(int)); cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + cudaMemset(dev_bools, 0, n * sizeof(int)); StreamCompaction::Common::kernMapToBoolean << <1, n >> > (n, dev_bools, dev_idata); // scan - int t = ilog2ceil(n) - 1; - kernScan << <1, n >> > (n, dev_indices, dev_bools, t); + //kernScan << <1, n >> > (n, dev_indices, dev_bools, t); + // up-sweep + int* temp = new int[n]; + temp = dev_bools; + for (int i = 0; i <= t; i++) { + int offset = 1 << (i + 1); + int numBlocks = (n + offset - 1) / offset; + dim3 fullBlocksPerGrid(numBlocks); + kernUpSweep << > > (n, temp, temp, i); + } + + // set last element to 0 + cudaMemset(temp + n - 1, 0, sizeof(int)); + // down-sweep + for (int i = t; i >= 0; i--) { + int offset = 1 << (i + 1); + int numBlocks = (n + offset - 1) / offset; + dim3 fullBlocksPerGrid(numBlocks); + kernDownSweep << > > (n, temp, temp, i); + } + dev_indices = temp; // scatter StreamCompaction::Common::kernScatter << <1, n >> > (n, dev_odata, dev_idata, dev_bools, dev_indices); cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost); @@ -248,9 +261,10 @@ namespace StreamCompaction { cudaFree(dev_odata); delete[] bools; delete[] indices; + delete[] temp; - timer().endGpuTimer(); + //timer().endGpuTimer(); return count; } } From 2f3e89e5e928a842bc42c4676a168007a8eda571 Mon Sep 17 00:00:00 2001 From: lyifa Date: Wed, 18 Sep 2024 02:49:20 -0400 Subject: [PATCH 10/14] block size chart --- img/blocksize.png | Bin 0 -> 19699 bytes 1 file changed, 0 insertions(+), 0 deletions(-) create mode 100644 img/blocksize.png diff --git a/img/blocksize.png b/img/blocksize.png new file mode 100644 index 0000000000000000000000000000000000000000..9dcb383f1aaa1df8f8c3d18e835148ad9b6ca27a GIT binary patch literal 19699 zcmce;XH-*N*ESkJKtKc)L_k1PP`XId&_P6c?~@@;Ef$pxJoDAq~qCxkIxevULx<`=;5~+$jke{8pTR_Hz%i#H~r2f9XV}r6NXEP z8u3B6vUxla@7r$_gqtikGMKA|m%|a+UYogoM{P#chmg4dOc0(K7{hlb*$@)o@67v9 zauA4Kg!Co|BsXvw7$iveo)vKC`=|fkZq$FdR8dsKUGFe@cklZRt(5t;fb?`Xc<3nR7ER=fO}Cva=tCg;F;&f6I7t6Kyt`g4(4{ zSMsr$zVka0gQku5zW%voy2or5wFrtgTPG=RZ=aVsoU$=-w?jzn4Tunn82b_KI)(K! zO&$z*O^yYL5C*e z#ah|>DPd;rZ4*+&aw!3f$tx0Xdr?K0}$da6Iv-}W3y@T>fYZ3tu!cZ@+AG$ z?R2-0;VfypffQlvX3JIs-ft<<=uBgm_mFJAv^cBJ_Q_cRfV~ zA#F|~WzRZ^g#B?XduZIHs33E%xom~x$jTcHS?)ctY<9Cv9!2h$cI-KE>wLX-&0Sj7 zf8%&I2zJoCA18qFcs_A*0Opp4oPHT6oQ!7^U;T__b7o;-8E6l}4-}{*4CKg%4{ZPL z9t4b8+Uj!}+}mh5C*nn7o#$Kqm=C;`hXzZwFo=-nK`BR=5Da^A76;q7d+;7>00-11 z|9Q+~bff!`gd?n^$DFgjavV3q2&-hkrzJGaYpEsi_m-uDp=I^9O&4p_nCR?R^(Sf( zO6BcSOUZxUZLY-eOsM@hZyM&DjPRd1JK1xU+iiqKcj?PPv>P$O0h3_2UJ35Stbonh za0SA~nNUvP*~Xp} z_iB9kBkEHS4102Q?-W;U@3MSeqXY8BiM;{$rMtH?`|9%5gpo18tSxuwqOoyb6iZ?*WUdDulwfLy<`gH)_|Miyno$zOw$3&iIt8N+RXF zh-Ur`OGdL*Cn=s!CMT)L9DQuNI#%ezN?6;+bw!=q{y{}D28A+qV`PjxDlOZXcp1f{hitpyFu0egW?_-(DorzqO17b&x$*JY&yA9$?jZN~I{hC4@LO$1 zSsQt}$FGxHj*11&@ag3%&4%SZ$cDF*t=K>Qva7X5cRZ6%a`{!AWo&*xBv89Gc^pF5 zm;Ym59^?=42??U)so&A&oKFxT&le}E4MjJ&9!7H-9Yr8Q2&$;aR=D8o(rMz#Zk*!W zi_<$FX1F;*vYVr}PL{KWO?w1}rH{=0EuemIs26rXcw%DW(5ThaaTW|Uc;U=$v=es^ zvl|F~YmWPL`{`qwW>furHFJw#>R)b|?2+bpXPnpSmXLZo^z5|8Mt5|v`sicJIyT*1 zbE{Zv*!L^d4enPi`9fAZW{pJ}=~&m)L3h8@r(Q)tnAKy|&bCT(rA`G|_*s=oNSmCq zrI$xr_N<73!TwY2ikGrBKdOh7NVJaxjJ4i^9hjr>fHT{1Q5ok!wEB|->VtNhWiE%+ zDzP9I1W}PtVQ;fjXer$(ezQdg=VKn631?v(I9vL$-8g6v2)BT`p5RIyZhu}QW3jpx*49I6cDM3Ktm^e3DvEw&U18|wI1c*Ryd!v5^iQ52KM*k`dGYXI!f6JY!gMcf zuhDS>-z&6+b%}WiEMPTFzFg68mQ^MtN8m{NdomA%;M<{!S0w$#bdNNmmfJ}tTseb7 z+FAMYN^IqT&@Oh+at03SjSvueZq{&Mb~EhHdD#I+en}qpS&b`okU4?yk%J;^Ne6xG zS$PAR5o|RECkY7wHXY5hAO;-V#{KS!S(Y;)^gQKR;=;*Lk8A2^z;j5^W`Okp;;~$Q2h^VE&N}X$T$*}Q` zJF0kN&f^m`#wpY5Bu-wRXCGJ>!&=IUNRHsCSfLT)yFqVSG!9ScJuso*y=VKJ3xY=V z0WvI+3vT&zWVA0Prl`+jHv0^co&7jXTPkHnJcFzpn<=^|&LAa{zMI~M3{_%r<;OjF zLf3;npY-Q9P$f9CA#CM_tvfN;-jpNL=?)$&Yx0~}GPV3L=~52CI0gGQ1}wa#d`ztNxGhifK^7KndT;-C5gn38 zujmz$JUrSA2qT!gC6DmC{X+KqdLAd}6-5l&IcIZ&QUAvl^CJy(SxetmtnQIsJ7LK* zOUrCivHkRAawGO+NV&Su+GOMS4)&0zp>)c<_U`{pmY zL9DuYvOW9EW@MjYZN@8k6=vjM>^Nv&aolrFd)sHY2#MoM)@=p+kM}T?QM^&DYWF9r z`;@D9Q3fQr{8Di|9jiEhxuB89QdvK&m7dt3!?i#cH4*taSorh?xf6Y=GN2s7AgAoeP1WV5)W`;h71vvB|R^o#Dc;fGF8_Pxhgs< ziK{1qgR+0!^#%&!PL;MA?=&kRYGASbkOx2{obl(#1Kds!uw+liIhuoifA+L^KRP6U z1p45_{`#rXoMXhh-wMa*=ZrC&o|J3&gDN1Bo9+eg;>h47O)chl0q=!CAea+Q&HKF) z<#;TU_7+ zfIm}%2Upw$_OuYm=J=04hqYJrJM*z zrWW|7xemUs2Qw)Ew~*FyG8m`2M;)(&K8*?*RH@x6F0r`<_I7xliC#_X=47Wbl zf)Xq*D^nE;Sg-9-heBZ%?Ep&L_jEu!b2AseKt?L7WfRDol24X?VDQ z97GH{QnT7#=v=j>{zTnv-QTA(MIZrv(@-Iu1fbwJ$5i)wTD5b+1qiEP;x~%^Wcrdl z||wXE3!WIR6`?l_nrg}j4E z+K>N4+pUb``UGtRobIo3r>YmBsPK0^lOrbj0Z8ZgJodID1ue^~EI%xBz2{Qjr&DwE z$puc1(>k)K)a?`>FVg5Yb_a&z(>!S~irGS(2gBQ|s)l5kIv20|pre_D8v4D>eq0K# z32NRqzJx0g)0}-&^4PlTQj2-dEq)pW)Y?S3D=ssZZ>@G8B1t58D`ti{3le>-XQ6pf zh$^+WK7YP{Jd{}_;)3Cxu0|`)YH>Z1Tdg*jI46twI$O2pZ4eO1Rqsn^Azfh53mlEq zah=g9Qz)Nt)xwQSZ*BXNmds_`=UXd;NE3lzK1FDTq@9}WD>3q$)xG){u*Zy?yIE13 z)UeP5lUkU!o2tuM=@+bVn}b*QC&HXj?>F;r$;(l1hR6vRn#%ZYnzOz*Xbh0edh>^> zh5uM$y@632rXGjLq9h+mw)F2sqEPk~%QWa@wN43R<=;~q?8wEU5GaDd1&Rm*OFy9g zxVb~YwSHmg=BdACSdZrK^3u}Yr(NosKF!BH3FXH)irJ$NX0!2(4Hm-HmA|_q?Z1_f zSEdmgh^`=z>4U>iSP=ejKhk^qN3~FkRVrFzB{-*iLe0i~pFGRCx&40F5&*#EPu+t~ zq9^tG-zTj`>mjKN&v7&c5Nz`$$mWctygCHyB<$7Pvs7ML2_w;kI`Le-WUZ`Z;~9ns zUrZ<=VE@Pr(J%?B3;8O3==#LIwQw%amt-1xli%1X?kUREnXOUqVRied(@+2QnxK=g zh5ShdiL<7-IMa%-WixqgWZey9Ih6Mpy+I0%6GJ(xR+^p~9qGi$Gx$m({+d=#*}BnW zSxRhjz}g{|OoiC{=V6=NjE(6d1S>iUDP;jBDCL^k z^p+vDdM0>+2tiCvFcxhWo|ogfqVP*>4CEw@&I0R)&*b^1=H0ox%s$S_g zU00bZJ#wyRj8;RwBgraNwF;%Ly+1|HIa3`yGX&{@T$s@Rg=3A(8qK&n4>*(Mpxwy#@s!N^1YqeK?X2T4t@P9oKDbSF-zv$#vfP*lJ zQywYA(Y-O!kz+cl^(Hga&$(>2nY6&K9^yCUqti9X=WR0LU#6?x!TT4s*qvpHs2a51Q3jFY zO4m^gtAt=%jAkU=jRwpkmSl!+=8hwE6x)pkQ?584Al;wdZ(=*DHFP#+2!U8PxXq<` zw5%o3YI47B@e9w}wDK?$n5%_$=DhJ}^lI*(TVR9!3I98jE#%Uski6P_C8=4K{+pCj z%VcZA*?BbUHZq1?J%I$|&}K}Q6fl`Ac~xb$R-EPho;AmzPCvn|(;RUo)`YguTCCZn zYOi46o^|A&b?uJrx43DzMM2WoWKWE6W#Yb7!)|EGhmx^;^E;vft!+|or}3|+?wgY@ z=IA4BS3Fzj4{lQqS5mm+d-DGb?JWuVp7hU&&>Une|G^|dfGPAzJ-6X}ctX+Kh=nbk zLuNyZ|73q61wFK^AnD&=-|x}fn1S(G6i;5G!q>!^wmXh{*bEMP%$uzj3 zb#j1Q0?vD|Ne3wi^pNu6|In0cH2`_{^7+N356mPqAkZDV|93Ygmq;fsNEm?iagBZQ z=d#y`JZ0+UTuXYtV9S6DxMHD}E4RaB0*lnU|xlPiU z(?~T2qXQ1rQdaxdHo!p-kSVfpr79@pjeBL9-~PfMe8G0LT)*%zz~`JOjcC>F$Sbl% zh0GYn%9du@t9vuh#k_kPCu=e2z|J^XSVYG%_)B<>VzfzQ*-N6ZE^wQBh9K>rMwFR0={Tq1L(M2LNC~9dRbCK$!)TvuA$Vn z@_}P=sQ43UvEm`~9fvU=^vSn5lR=ylAsWx+^5n}u@109<8YuCD9=c!9+&lNIU}UmZ zl{2#rf*Y8}+vWlW%X21)Vxd7sdA)(l#CImG4W0TD%bYRG@ZY6QW=ycS5;!A(=Hk-h zihp-g2>ef-@I2w!Mg}!KyeubK{+KL@qQv=*zg(q`Q)*dh+`O{>oPdqx4Xmf0EzS!G9?8 z8Vw@{$(>y`-nheWRob%o^V$R!=gT#4pjf_K+G+N1M6H_hg)_qQh2CAr?)`Y9=^!5&aP^K z|35?M(lGuCMsKqo-5S7H;ov0w-k=|E;M)@LC(C4`dNkfk=fj}E^p%YY&n*8MagP#kV>ws0kuhsn{on-j0 zmD(W)hxJ@FetX#c#|ZuFgPT?H(20(MQkwxuplMm1uA=${KplFr_-SNY6v%@rtd^?v zsJw%^`Z-iARF^v))#rdN1Ij9p^fw|373ri!MWyC^(5{E)(gBCw4JGR3;u2{WK2e{h zI9Xoq_>&Zr)8^e93O!vdPG9Ozj-;rtZ1bN!&wlggV^P4-W`j9&!oY58zHNw|x2Jw@ z&4tf!7=PUkFH__OIEXaRlJPQ)anEfpbwgB6-%>JU(`EW4kS#Uz*P8d^2AAyP+-f;- z#e?T-=R`0A^+lk+IkmOlL|@YyDrWQZib&w`YaruN6|o=6@Fx6@PbiuoH8M_|3pAf(#a$Nvw@u+4)<8_+C#5joRKN3n@X4s2|4FpmxRHLk@a+o`O?j?`m zWv#R0U)BW!E=x*MPLtO31j>^cy{VZp*s{WV33SKpQib5!u#Il=So7w`0>_fGc`aZo zP?MBL&s=!dMs^z9dRxt0l@b*CfM-XW5_NxMkq)nEJ ze-x1LsfQGF=PGcqs`jZK;iUtN)+FS%#fA|XDAe$dRcByLOK;#}7V=HSGSlOWIdE0# z2ou{I1lliSOl&9J7QX@eU?way;#&)MQIXqmFW^~)oWK?~yl2H(2_H(~kV%7o)jTi# z$BCS;)x=vLP{v@LgUSX>203vP6aQYPuZg!swt~@76MOCNrrYC6SpkrAlZ3{Ep)^?X zu^6=bJhnt!;}vjnMJaSBlw=xc6d56aL1-aNmoLV8`4a!%{TWGGAO6W^pzA)3f4a@R zZ2Vz=JGxYD)%@`F*nOu}zDvMyeu@8CnfoK`ntW*~>!}oE@uEF&@h>bmAip`sJ(1RA zJXwSch+!x6HQMpLY8P|8z4>nXTXYp*YaVd12JuM)*crfiTJ%-8Fd$T8SS$>AbHSk+4388~#3@P8!s+KHW}B*JEX&@>aL7D{f$`O(*}=sTJ|MBB zl7cTnPN}GPrM8uIn(w_CGxzrGs_rvsd}R@SvSXtnNs0OwzL=3(AYHXAX3USYyUgDh zbOXF^mdS`T5-8jB5-w{fK)2ozA-9dIL1le%aU|kr@aL zG55;G&mqTd;cW$)DcM)lXC5(d8)Yhd6@1L5?L284S8&JyMRTmXJ*Js=oyN3^(PEbs zyW#+6)#Y*Q@W)j)4gm1>D-nMq`}o~SpHcfu!%EYo%8#mhf9m|vu1r%+`J;ij`FCuf ztqc592;iS7lZ9%htNP`#F9vm442lWn^S4XFZdL?5Kuz6kc0)G;$Nd9kwzKPAJ|v(S zBNehq#5o>;CfNYUwnqEw6lAl*QhVwl2Ku*P44Ih5K7JFl)uEw8s%%VEEoX6q%)}{! zw$<3ztBXv?>70FJ2H5LIzDmAW-r6$(J^KSq3$mR*!2)^;f8nkucV$ELhG4)GM-&MV zL=#1=0D*XLAX^r*k3$Vp@MW4b0Sat-T+q4=v8+T6E|dFH#bQR*`c$J-%;?ICnYSwv zQU2V=vi@>X1I|Yqd?GCo2tuqIk9Tr?q3j3m< zHy7ER7e)`O=Vi?EoNydzsq(v}O_OX|;IuX#Jr;&ef}&Yn`wbV`Dr3Yykbp+AR9XHB zyWWj})WGf_?l$XiE(5671OuI1(O9QG)&ng1h|4AAvu}fFnxmT8mw3}Gih&Zbk7?W@ zFWoZPoD-|v7Kffo3OY}pQm~diAFNiR^Q3z>*D$#?al6F(iqna|w}0gcVf-5aqW)Q? z23PbzqLHPeAVBl&7vw$2cpYZa;F8+N%{v+u6XO%h8*!h&_9oam^Tae;U+rs2zm`2A zh*0&eIoJVa(pYIbR|2Mejk;B~w^dPBY2w`Q4Y>q&WybqwZh-aMe!g@i)QHi7LJxJv zLL;iw=0dVfW*1? z3l}yg{A!<5TON72vq^hyXNH$@;#1t=%8oIR4#?T-j1`QR8&&@^kTlU7EvB`ik^!%K zSvde4J7V%xbaj7yT7=ISIEuIr=+B+rteu92yb%xlqXi(?ocf23W6Yus@w9rMI>Q{g zhTZ2ASmFw4r0@W${pbC-eh9Jd=|U)3ooRJnF8M_@$OWXy?GbUCed5(g&XlqkD>;;aFDhr)+SLy-3~nJ?nhtxPd6$(bOuY;YyKPMECR(TFD7ftKCw&BOq0$0 znePQog5Z7QF(jDDof~igaP${$HraxAQXCfcFMw&*;lc#_?_^Dp65q)^};a=>o%a4L} zqT?wNd<$wqp?^Os7bQ0f7z46qQ@w4pM!NX4u*KIueP<}9w0^w%FttH&a{)$kZ|zLU z1g-ZFKwQ^-M}7Uxf$r{qevv{X;M41e!2N-1df9aXig5AYPnR=n*aOyuCFhIsfssKu zDEi-DWlJOCzm02pH9RxBK+Nei@sD2$21S5E`*iC<4XC|6IzCb3!Y2`j~~Ct67+#n82rp8K^R3RP{D|s z;Z&*5kJ3~C$t^u}>=+|YbZ=g|DVgi=NN2ebL2%VGFaVB$DhD9hG(UQ__T{PdQe@oy zK`w%x4NTDQj`jR}hM$K6fWr)c&8s0X_Z3={MQmFuD{mDWH7MnB-u>b)=^8%7e>T}h zwHKW`VRse=*2Bb8TO)2KIjw>b?un@T+_=)D{wEP#X(kb3aHm?hdSW_Sp3i~c=p9pP zw=N0FQ3s9!wve+OW&vaQr`QymGlREb{BIgN1NX9!O%>%7+=EHd-)u=heOkbqYplEL zx)N1^f}@MaT&G4If{z`XUbk}dTQuZ$LdZd(R}IEOxrSxWTc2*V?hSbKu-OZpoehW# zAD4ra%F7%dBiK?|&Z<)9#j1e}bt&&J3C%;s%(ute(XdrVw*AH=AexT?(cEq*!%GU6>i~O|tk`$6mVMsWT?^|7$GuJ$R8GNSL(83ifc2^O^TNP)2dFQP=RX@mb@v zdup6&UUrU~0EC<|15EwGr8Xc>$@X=^Uy0_Xs)~+feVg;7FS)=VY4!91K&1+CX@J}O z5BjJLFY$k4vS`#s3F-ejo!E8?uuufLP*h3K1PZU}!x%|uHdZ$VW3XQ+b;B`;Bde~~ zmWU!CG_EF!u3ljg{q>&br&aE~jS{EFDCy5EXT*b+>vo4_XW^J0p%S&>YR0ec<6Xp| zj%_FNid0_Kd%^Dh4KwRys%=9Ysq>Fzlg${8^VUegKb_gN(ttdmAem8@67-Oci`Ez@ z2}>R%p}Zacj|nqo7G6d0#b)HfQN3IBW38cB_A{B|LQ~0DzP|qeCznU6lj5l@+>=u~ zme!ME)%ftp1nza;_mE-`ROS3#D&8l2u9~mxDo9%fI3Lcml{yrZ1AwRmWT&M!^U*(1 zDQUNjrdB0$f2x~i6$3{eg5t|Rl66{6;epkBa zKtZ9qz$vj9$xC{(ud}?Ibs)mG_@`}Yxhr~pJ?YNiw3>@ zU1D5Mb~hZW@P|QE_0?d!&l!B~Bex_UL158oUD(O8dV-?>el>p$@Ue)3a6y&1fVCvO zU#CL^s=Y~ks=ShghuWlgN*dT^c?{kJ=|W;1MyLhBK_+{|!}sq>im}LP!?H)gnWPDa zg~C>Ar21`kMJUUEx+}a-!sn2?YgXydsqj>%e7rhaJt7}C#vhL-FTKzn^yYr4@BFzc z}GhK){M`OV)#`WK)Ut*5P68$CWVeKt@|>(~3MK+?eD zZD)nx5wJfXK$Y4o#BKcSrl@I*aSakPhE*K_W~>_r^cK{v=XPO#m4`vq1Z&C~QOBAHVAS@;><^jnbSXZis#+{LJmPj_dL|8~{O?KXc>3Ai0x=D%8Lp)pebjc;&`!VAJqK zo(dHDGVwn7wHoc->o5&$Tz>h6lu5T`shGRJ{jCN=8yhJ+3Fs^0OTnFxRs&;_zEnE{ zlyR}uHl84ad~eBHeCt9;rllR-a|90RNMKTA)gu^luB)3Fx(_4uevN_N#B@4c5Psfl zTq@w;6ly-XLT#qpp}SYOZ@6xBc>dR&cMrYA!2w(L;d&wn++(pl6a&Xs05FU|zg`S( zgn6}V1b#I~AxZESK|MBFVp*YsG_)%JKF^2Hh#{8?3bq2Rw^#=xkO^5P(j5QvT5uH!`INAI#D%&RU+buB1KXdJi(4~ZDMB84P7+oNiO6@U?yL&*{V2u1@(O5W|FW^1`t#xH zdV3~b2ZEn^Olg$}768nL1Bx&=^7<}_lJ67fCGmO%UuHL-vl10KR~k%7Sa_c7`A(E{bdR90?9p)>mXIVlX4&p84Wab zHuGBrw8oRWRwTlBD_=ds_$XdQh@ji^H^9YvR(|4uC_48#Cn1Dyt?S5Kavl4|rvxHaE{zqDQ*QQzU1L5yF zcyGK@rGC3@ASS88`H^feKNO8%)^V-gyY!J;q;N=Ocu&4at!StHk5W#Ga$Z-Ks$Nyl>mTd9@{KI{>-1-fSENl-q=Wx~3 ze+zhSo!wzH7}_$kJ=LldOZRPrj8K>GX;m7_K<|~&2-SEfmo2F=TxR?#+JlPlec%#L z-J&1QQNxnHHNQ0Yme-+k%;all>Q|i$L*9ojU4Vf)IE8^A;>^j598C0LA zp6;?@yV?Wl!@fd@j{?J+-4-ZU&$EN@GG5(_=>zXl`|A#wl})}md^TzJES#>hjSeF4 zG6WsHeaA-mkH+o~lveuGw1;KGrH=khW>qNJU5gjl5II8>g+TC41yJb>^SV*BrRzm5 zhoT{!yhym-)3UBVCO|~yo$-JWFp%|%SO59kJWw7~w@CxAeBUkhGufF)lk5SO!tQwc z_o14@P{z~MQ0VdRNPIcz=Xrm`D4+#wpVWv-hX)+y*af1DE4;QGOq3V%7{pa=aar7w z%w_1Vzv-<7^RcP$Zy{a#>AKEUx8m38%wWxX1WBI29=0j;*SiT-~b>Mwj@9HapY~!52-Ri{$MM%p#4$enw;gw8E#2! zd`MPBgME+V>=*d=`MEE1ncURUbZy4Y?ZHNBkMJk)1%7_Et;v(60A!MNler1)XH>1! z_tjE`F7rI((Iyuer6h5-H}LKl3q8hVIZ$ah5ypZlsGeRqIt1m6(bq^ZS$OLh9_6S4 zG%CU1?U(6vm&e`qsMM?@*riHTS6dUH7+pkp4eq|<#&i|V$5BB6G zHH01tUnm-)k+5rCIG2>&6B6f(9NF%F$&vqFSbYByKLB<9kCf)FeB7ZC`nH1T&}sl3 z?2-VqmgjaUW0@GXk}CW8pNR=t@hKi8xquFzxcJOT8AH*2quShCFixw|6-8Is-~~B3 z3*};2sCwEogsAUpuydUGAcy>SH;^#OS-cLh&~YyoxJ>$)`$(T=>}s{yRnW)|5d0$i zHZm#p5Lz&GvK^R=0X9oT%&t2&)_KR#$?5UI0Vi20C?^_-Rg9|kyN);at#tJ=Q%g!| z&xG|ia&vP#eX7Yox)PqBqznc6b0EIVt+1?M9+zFS>SJ?jOl- z&s<1{sj?lrQy}fw!3gM5U<>pPAoy)Wl!r*x%~*yA)5F+8hay7By>-)s)I{=e*Tv;W z#%6WX zll!zdET^!!h&{v5xncd|1BXPj4(fM<-AiKwr_sBY!b^gVq2L|4WDe>I%)EnQ^f0DM zAd1S>zu^F^jWeaML3S0L<2(8cN?P>Iq}FukQ*z+CV28p=|*7bps}Uke>a@Hc*sS1bA+5ndv3=Z`=gMtm7&8l|2iN=e%)-e>ocMsgV9Z7 zt~IVIZFE>xKi<@eS|dDR)bGjVmy8|YC&mmOzrKTfS8TbA^xN?2J;=59oB{qzL7&8OP!$3PC;SogT3NZa#d}~mt5u;qr}FnyxhVL9 ztRORpB{In-E0`e3->8{LNb$`I=g3!$R^&`St{$U*aAIT*qGgZkJP}Blskv7}?!m<5 z3fD&g>ESt8-FsHK0=sSOo)JMIY%$EV!5PnBS9^BuBCnC^{yPqiv*ATSFWxXEPf}(Kz3^5{jiQJ})YTDWzURFQ%7Kf_;qu1dS^1mj z56w`YAjp&qX>g`O8eEvcSXA?G)km?S;XMJ!b&`;lC(kzal4`}6hjp~Qvx81wjx*&5 zIRS%k&{?sxIR^(7Kg{a11pE^^v@`Wj)t~?I9K~07V3VCs-H^P^LQ7kqTYJ zeeAN5QT}Vj88Cn4$!Be9fR6jTuT&wFy%4U=luYbY9R1Dv@brruej6cm_GopNEJ-0( z;l1hpzVEG&^Qsc~!;`LOa&se2nIlZq3UU}Ku|sf%^{mKnUYDYN=9S>m$t$Uh#Cx1h z#BU1Eq%C#An=ObkNR5V$lk1NhZ(T}ZCo7>k!!+zMkrHat%|TUr1M}YF*Tbl z;@j26rGGtx;TcCf3pyuS6)(?gq+?{W1Cd?frKY{%G%@y-g(M=ZpPh=1WHZ%X=Qn4W zCc__Yfve|lwSS#C>g2VrvN7c3-OCtzA7iyChm=7^6s;P$u_vCS&B2R%8*aUKWimj0Uc+^p zUE?DG+*D^XB)92i1w=gL?OOjs7m#1!3Q)A~mEtH^e`x5CGV+y$r5x7ubX#A7lAU(f z&?ug7eV|@W{|!)yU@_L}9#O8atQCEoF_xNbU}H-W$3sscN};kiNlI-24$vaQ2kAiB zTmiMIs8$&NwDJ<*L`tcXXPoPiB?{5)_o0sLw9h7TW>o<_K1QMY9Yw-2PW6i9o;u(p zRbbk7)cQiz9VRg|mieZU6)CLbk2?=;$=?1}yZ(sxn5iE7&yqW)qHCr*mWL>OBG|k zn7G46YCY`p*|rMkbG5~sf-NTR9M$~CQA^P7U3va!rIOA~e?nQbHg&43Kom#pnfwCa zv1CMt|5Sf=UnGkb>9EonyBo#zvYWTZDk1(g={b~80flBhy7??G>rwQOG` zqi6C&aVL$DscSIf;~VKd_el*`|5D!wUeTsWq;Sg)v0wAlOta&R#|_k%Y||S&aH1O9 zaENc|8rY>3h|}WSE9PVeCeb36;vw#<|2d7S5|e3Ai2Z(KUOUZ%Lhj8=LA7v{qW;sC zryY)gvE2L`E z>!@IDWZ{0uus&LaD>^f>#Q0K0Um+gmj*$uRCn5SnHj|d5+!Lty0IFK#igV--092Z8)-y3hb(7Fhs8%^E!6o)w-KHmR z#XB%UB!oj=$KrKgFqhtoiHIbQ0FP{Dys1WSi>V}_gnL%bO*&vJrflYCjophUE>Jiq zbF2lRO+aVqSpbZDcmhyRV425llUY^O&28o(>qRBL`tzYdB1dAnG z0F$%wT`L22jb{}!KU%8QG;nz$QIfkZe_jGrH{SkfK+!`0+->$;#=v!1*HgBYQPf+b zm9ySyu*QfQ!4NN7Ys)n^C*)#2)5H(3yid(=0RN(zQ#}G zpC)ZC98=5U-5PJ;8~@6uEgxzQV3IemLsVj`d1`IIt`RJEIP|A&@a-R9JUIN*8mx0P zi=ES8=E9`ZD)1OP`O};S|K;67=@UDANeyjBT79f-9&|?)xlpx~_%;JoLZRJTGW^SV z9=yRPs|y90`x652ifGatK_YJJUfD`mXJVBIC4a=H-stOgYhPad@Ld%|(+98SnJdi^ zmSO>FB9>1PJJn~3d(Zkrj1R{~Chc&}sM3 z>dLyG)7;~&a?j$c8R{w$cFe0}f5jN?tv`j+|2@-Y_wuO@U#+PP6!zZ8P`m)c=kLDK zM(-1n7d3QQ@M1T-z9RVg)0e2pjTga2d3}I9X)+fJ=(jZ_H7WR{>J`(^n1F}Yin*lu zn)D;dMQ+aBu|Kb7|5PYG#;NyaUsX*}5oZ+*jLkQ_wGxr)gEQX<*dA;wfNK~ zy9 z5IFpiOJ9_^?QEDo-nLRmQ^PU0lqXo`zN4Vv!j8^p3QOie;Rbtf?c^u7t52NTv@g#l zKAf{T`nwL#;QafA^7`Q+-za<|hCG7N;z^rW0&dO4EMmWL-u+v0scwp%QN~YhUQr(P z^EXywfk$t2% zf^CZ-C6r3gL~wqK>8&&T4MR~aW2DJDYG;=C>P_B3vw~=EoSHGyUt(u>@5s}Jv%DQQ zBdIje3tnY2OW)>w+ZQmioWR}Fyf`?)oZ-(IdM6rgF0tYCV@NWE9dh3OmzVmRuEF># zaS4o(bWz$XqvS8THKP&JjDd`4?dRHELn2FC_aLi&kS>{!8Re}Vv7&RNY0DN<8idYk z-mPY_Y+vdo*f$YQO+a7fqjmyY=g$J$?#_40LCBS~S_?#LIGqm=ILgZP)2^9u{yHD? z>~Bm*JVCj;2ZSQkcUbCy;uHV90iaaGws?{zqEVqIvGj+E*VnnpGjie-#Cd}`wYti! zhOCv5{S57uU0t%t^`#6M10^lly!t%`%V6;*jQ)%( z6yma2|4oKKCb3Az!9xU8%hB@OpSTN)l|9edy)WgUXs*-4&r7&tk|qUE3|0 zgqE&kK}jV01(gY~Z$9A5unLEw&H3Sp6v)-`t8MMCjCK0^wJ5kZQaH$k5i&+P$5SSW z((g4$x;VPUT}`kItZCJ~vY+K2kTyLKp4~A&RJlT(M5MZ6JuIKGg!3(QQ%>mT)h{d5 z?kUITEI1caI0m=(7q*yHx~>@>ZbkO_Zfd1OY#Bapx_&`H>a?Vhy~i%H$g~Vm+>Xff z%pb`Y85aFuy7(3eDDX~KKcuM8tV9ZAhdXH29`oTj0+k+ULwH{Bb6?7g`-whu5z zyc55gp}F>IPO&0TgFGqv8ih}+5<}*%JIS$4i?Jfb8JIxCUD zey@;@i7WF4%Ul2XLK03HPRv?=4N!9C?z_46QXs{Ls>DX1?4dp33{NiG23K(9)! z*LatVF;0ae0`^yaR|CH$K_3NhZ7Vt9W4UFyxtxR}psZ*8zpwB9x~FJXb5-{2H9*02 zgpyzZ!Ta(5RJu$<((eV7(G@FBjN~eghh+QAI1ZcnjtKd zQ9%=XD-#rp9Fc0GPJ{y5McHy;ZG{ znZUpIIE0lsg1()M#s$Bc4>=DCnTXPbK-{&yTn=czC=t3)?#2Pdaqpv6U{~N#3q~wm z!hOiN__pKj9PsVa^cmv(=zKs~_EZPI6(ce4;N;W9-U|>(x@_;PsEaE&u3dL3!znDBJ)5yk!+Ufh>>({61Pst+Y zvpxOYeby`UtuDT8BZTrLGihcJ48c?jz^P3U2~XcK>IiRB9=}jd!Q*)`983QDEX;fC#1* z^Ix@@`wi;7N9MlY0(*p9{oEuKkULIJ1y<9)KHmd8$q$>_AcGms%!7AzlRsHn z`RDF$pc`E*pLnE+yQXD&4YawRBHBCn600uPT+}lx!wW zDhQZ2jm%R9Eofvs8LbxA(9pnkkkNg^1@1&c_RITqkzRVGTR%!rd3@_9WnYi%rfQGY zCi;mGi%{w)ULy-=;&h8|b!D7?RlofIn9hS^%Kq7YPBM4#5Xu}gq*+gd0Z-Mk$lu5QOoL&0A znmN;_rp_pg)1s|I97I9V3N+IyrIoU%;6N~mQU!H1fm)C@B!NNeQn8i2Y{4i_1qoKF zY!WOYK_P4@NJz;Qmx)avR74064GI`<0>MBaWV^R-obgM)wjcTrymQ{|yyxEY{2u&J zy{bW>tc(;lO?VKqJ8eJL%eNdMo8$}kw%2epNf&3AXgjJ3ywy5*cf2_3t!zVU1u%=k zntz!+aO(C2){>tKl=2?VEEf+LR1dI>eJF40vhTFOp`tCy2KMNcQ*%9>!Y*>at8XUU z^@>hLXHPPqBMy^;o-Y&b+iw7Bp2fwlp4D*N41$&3Q@+zN^906}<=F-Nq7Su0?vybY z89uTJ59~@PI3XQcvj9$5C|kL9@aD1_uuiUNf2fHoufU^FkCI!JR>3RoB59Gkd&A*Q zYFNXs_XFi=kxI>6*dzevt zv!=!Kfug#;YjF^qo0`i~j29WWirh|~IqJkUJ+~Q~JmXvven0FBl3R(!xNciUi;>GT zd1dX@afu6`&-1ITy&Kk?zmO_P6)W(W8$KA2Z+WK2pUEskh zFMoG%NoLllM-PD82-3FKP9S6XkD#D`&ljzYw>^)az3={GUNCT=>D9xJ`RWNL74kxX z&6C^_dTihI+CDmhO*Fk~GzH!7CfU(2)$Vp1S#B9aJZdRo2CI1@Iwwv&c=II-KBc7y9^zF78z)QpKzHaw7360?bE{V3hziR*3FRAKg_DlkPXZy%A+<*Hc)?Br z)My~Wdp{Sj9%-5(PLvzO6MuTygAi-3as~X>xXU{LHEj^my+<8jYsT z+lFOix6)L$&XU!%u*C*t4;#g#W{kw9d(ke=14dEio-uwJ zP8EV)V{}vOF^wgm1b_nN>)AnfJCa~zne7W8D)OOyMDUxrxorjgx-x{iL(*L8m;D?A zqm=T^o>l{7kwS9Q0dlbCbtE54Z(*Wikd-=w5D{H57e!?0q7wWyCw$Y6Gw+Rz*QSzm z=PB|jq?G`q&WoL(pMPx4`LjYyGAO9(hG_sh4D-m7fbO_$J}<5S`7^134tVTK?M(%D zi|VO=H6M7C;mu?oZ#NB;iVmSeF^$VHuCTqa$|0jbV=B|@#GqnQMbj%m2Sal<*ow(d$ z-CyPp_K@FC4C$tigKEa#memMNUpfZg=d{a5`6~j?Y047Vtxs9ih@RU=$(?iN11dg+ jkT+*5UH>0<7{v+PZ)tw%DW3twqurie Date: Wed, 18 Sep 2024 21:57:38 -0400 Subject: [PATCH 11/14] nsystem update --- img/b4numblock.png | Bin 0 -> 87101 bytes 1 file changed, 0 insertions(+), 0 deletions(-) create mode 100644 img/b4numblock.png diff --git a/img/b4numblock.png b/img/b4numblock.png new file mode 100644 index 0000000000000000000000000000000000000000..9af8923eaf4b2e3d2c776febc0e0896ec847bacc GIT binary patch literal 87101 zcmeEucUV))+HWj~ih>Oh5!e=#CRIehfJjq7dXW|^Akuq)011ksh=O#HD!oS$2rWcK zMFV&+_hH}fe9v>+pLacBm07c9&CGi1?|p}O3p1m=yN~V$ z004WhUeUJ#0CpJw09!xp+`%iMQ1`j;4qJk)j4lIe`;N`?F1CB&cq+RroKntRAix0@k1d1U}V(<9|9k%PI>lb&|XBhkR z`B=qkbnJW*te2L))dlP&+*cjt@2JVdQj#}4cgrhbIhLEXVgcbk15*tGV^>&p8Wq8aFPa3ph9s_Gjf>y(;gVjf!TcMg4$5vO(zO&(tmvjDNRnR+^1$sR$($=Zm;n zV%2(9dxM2RkB{f%1)G|hehEWyTr{FqUxT(j1P`qLL@#?|*0a{gLzph}d4s+>`k{=U6w8^Kgr0Iw zTWEA%7Aop`wd@yvH<;@*SB9lwD++HB4}tprc)KQqTJv6BtO}cNT#p-W->`*#(}vTD ziL~Kh-#qo!XXMgQgcPwxaSZ56KRT968x|qwDGoOdbI=+_Fb2Yw6V%SG?4^AsCT*eJ zJ;+i0oX_>G;^s OJgP5%P$%&X|UX_Ql4Uc7>6tD7tMcu?D2a37%jV;%OygjgWi| z;$w_t^w2EJYA}j&Wa}2+!PSTkAcmOK?oEpt^>vDNsqSPgB!DsY$T%?ZcMI1V0_obp zeZu)V?DAgPCt{K%G3TvL%kgEsmRvF>6L!;``<(rhH5Iu$U;x6_BdZxXZ}T>YhHB9c zg4sdiGj`NwW*4jir6|^IAy+7+3pc1;yQ$?*V2>cC^(5KP#mfe~c0asa_vX!Uqa=yp zlJM4f5!jN5PEi>%d?wv~x)Zk6Nk67kT$?_v1zX7sqAxr+2F18`3y3J7ZJS*tIIpy( zAI{K%CRxOpF>#?Q(r`%7tsey`>1Yrl{Y zgxbtYuUzVJP^(1kpV|?R@)|1KVcf!tO!~$tj1zV`e|&XU7I(US{EjW^KZUr9+7#IqL!L(->v`M670{Gcu%|@yXeXzp92!a zYnDQK?KkiyS{vL>Z=vY3QR_!uH};9->+GoQj$l~oTECr z4kWeN>O;TR@=r(G;VOgK-Rd>&{_F9+}+70LxE1U zmlynt9!?!SdklM56g#DI!5{hR)}DpfLdO9wi=hYW0st50GFlN=;gaje#KO2(u$Mh* zVRRmPq1G^$0(ZqdBL#|*C0tP?>^zvO>7P3SA0+ecaVtH_IHHUxl`*^@uSa*8t85pP z8DgsGXhSQ??9a*yMR>9%9`jl|DZf|K9`Iw_ClOK|a&QNgw&j{me2c@rKt|I)2sah^2k zyh!6(p%%IV)k<&bD5U<$vBlsde;(r&gPEQGDrXN%OtIO@wVm!(t(|`rj^PPS5}8-I zUP*z`-0w|C2YK9ll`QeQkWBgA0 z{rh!Ee--lN1dr2G|9tEFzxoDj|5Z4K&nK|4wJxO5pApBa0tm0GOfnU^T-QFi&S`8C zya2#lq9o^yWa`~>Ui_vFFKsM_K7Woo287sT{y3{7y-)?40G&G=V6isNRqyg|3(&*7 zw;l?XWgZlxlQtGa1n_S zzhP!^Na>0gwZ+BCQ4|%j=W2s@73tLg6pp>wI>BA&pl=xp(7ADFFOIYK z&mOc)CAh8?=3EH8jCoe8J`rBzp=)`Ntr!!?BxAm0vBJ872_v*d|1ZHXDxR^dO4B(> zfLv3|prR3Z7U&!s{A>elxdces`tR`o_K4gsjCoVGEvKC|xcnyR!HP;Jx_NM!T_I5S zdYQlOCq@X3-9XTSdL5zU4?3A}6gTJXTJVM`Xd@l-1)ct+tqp8>TBj@{_qjTBgER87 z(nk)XP!l-;A=Sh^pjd2NZfg(Qrb-1R_Vp(o;iEGzV&&rYVx<5een)VLSXful;IYKfs$^yPo1T#7<xix;z>(26 zda#ZrKuu9a1@H#E`X~i9BiU_4s8ZbmxZOc@)%>aWN85k`M~_3_LEimN2rhJl&$$e5 z5P9oR7Ab>-(7M}6sYIRkw10ZE*6I8AAM2gE1{Q~6M!0h_5!d%tAeO66xI!HYCD$r* zvS69%EjkY-{%~mvw>01K?!A7JC0AE#8-N`zA-Lst*FFni8e;b!;w$jwji&UA&N~0{ z=S&0b()-2h%~x1B&4U0X@c71kH=7NEcvoa==taD&$2hAYRqM)A#z$$&)Cyf0%81 z@M>;Qwps*^07;4>W)lK3Pz*^o;_}T1T+)TN%VQC&(zR`@cS#EofOlV%T)(pF1OZ13 zXSuro{bEreeX&od^_tDPe9{&lQsf(JE%n!O>Xtw?&YK~KuJA?WOpYDUL zi$-`zPy3vL8EJ$G{5%gW%o#@ zBkwTdy`-)@t`!41A-(&;rHQO?OKnGywo}X2%L6cG zp?vxApLWpJT4ouqV|g~{oT1mfN$q&SoH`^^6_Wdl!U_Weetg^oRLdKeRM6rc*zR8n zWMv|b>_E1lB5f4_Pj`d(e|^yc_0{%=miiV#F}k8m;uu~jPS=v~oS1h(5xe3X$#{Df zfY4=*pYsu1Koyj4U{6j(tx*(7WC>~ku5%IAPgf$(iP)Th5(2sf)XA}44P%^T|+ z`uN;Jy-P26wnde8niWXQCL7uIFEO%PwYin9D8{xMiWJm(5XwJ`!8ExmL>_s0)4$9T zm{jCQwY&+^apD?T03*NN$-#n?{#uL8<6+>ezbH*&lz3RCEPCr z^&!!sPta!h$(i&Z7QvjKQQU{ zUwVr`Kab|`dH^AdFTYPh%A@i)ak)YTs&~G%nvuge+?MdA!n~vYYvtSjcpQAI-qcB9gc4O1wHSh1qaG72h9eOK(iQ5wZ1m(lm#F#9=_<-cAZwRrNlR_jEA zfDjrkVt#=WZN+a{ar6BIM`&7)u%W(OgIyd3gy>f7&*MJcju%2bUI1r?RpMrvS?L54 zcgH+qJ+_;~TyH6H<%k|aO{t=}%iQ|bcF8ae=nVU;x+%0D0@cL(p%|Q*ZWpRuQb#p* zfnB50c^f(&m`SMNhF0a6(Y-cmGQ7uF-RW}MeEHR_;o#6JV5@Ck!BYqqzn0iZ)!ena zJS`8Q45{=jc+0czcLUUEHGA%`Tn2&NT zRyb}O78~$Av=I@<)s!6O^gAOc?Cgltrgq$vAIeg-k}eUJlR5c=*>XXa==m+blX=-m z*Ye>87uK*&xU>DIH?d?j{n7O8`24Q|`hHq0p&4MNVz-I*JtOMc_fCWMjYqzMx0%tb zkxWJOqdOZOzzs0#T`KGkfly8jD=vx_rLTSG{;Q*HE|h0acrs-ouLJ)k5ZwJY|F@707n#~5^)QTJsYvRx%?Jk0YN z$ZwN}OF+xyR_Ds5i{wzU%mocZCW@tX!XMRn@6j|{e%H4W-U%nf-(X3-s8P%bqY)J#{D@7uLMYHB^|Ti2m<2#m7a%Y9AYz9ci9 z%vv4s{ub*6d%Th36IaWm&5UudCcU{Wa1UJFQO6pNFekG8xLi<^Nbe-W!-C(@1=gE5vWId-ADtk=67M>7q%FKmyeCq zqP3TWX0~ZclNy=%&s*Dg78BR1UMV!KxKFovopyR}VPhJ#AM@*P%>X)GD5MIt#)#j6vpd!nsWdfid$chErk% zLdF9%!wRI9;jyvYN}coRPuZG_WfC=ht|&O0jz?efO2;X(QI~Indx8E1I|_9P&528u`*%6p0>zlPaS3u0!)h?@z0dO-0oDx5Ij;yGH180*S3j z9k~QW?ijTo5$@FDvcc;UC3?VDm*d+2qR<`jwe~G^&dY3#T5u8u+0Kb9(jowL6nf|r zm4Tl8!N+POP?6Cn0%;+0xXkh`0n&4prouB%1|W;mDkwqAI;-N*taH< zXh?pTN%SZ=%|9<8o4DkPJes-0?X;GG*0z)-{k~NL74w&5jFL;IAcA=L+45z$E&P~-tI!TSWluaXn zD0z};e$ln6or4RcbOQ96240q0AQsa4g0}N%YHW}Tm3g6;QN%{xvgz-$Xc8Q|Jw__{ zsWowPW37kCnf#$G3B?NUKu3dmB1m+W$_$g;42JUo5!AR}kbe}RM$x6^lq;8(-yhsC z@yh0s&2;S99R?0O6T_rkp6I446dy+Zfce7H&R;N960h31q8hZBFwH@*U(rH~iCz0q zQS+J&xf7C3tfy^*4hGiu>|E1-qBv3!9_4{Mk(t5Nj30z!NXFc=vA)6!Q8;bZ^|71et|pv zkkcF9uAG-2rvh1XpRTgjZ=WFGN8e;tsVmd@hz~_TuZer<@F%TT4Q57EIclcu1hI-P z)Ni-Xu!?>T$01|WybF`?oIq7lGn#}b0F%^$@lCki9)y!(@Wr@MkPGTKDLMevIUgp@ zmm8Mqzv`mMBQ?^{u}5zRQ=SNy9nr`e7hjO>qTnzJQ(T?Noo@Hzf)wx9LnOD?VCUDl zFP7Ip`QzhoEnQ-Viz~toSHx2Yq_A9L1nyiUVDQ= z59Dz?uHDcp|z! zHoFxBrPCu}0zW3;6`7IEuE0G^@N8Q%20q}`%HNRBX#o?NPY8kZY)-gSjlBhjf2K8y z{>%>HO>dz`au?F=aGlY)S=i$pE{Oz^1%j`E1ckdY1Di3Z$-Xcu1}=V$JC1D_sDU8R z@b;q0VU8DyJX?k%&#ybXcl*+zsyI;&(u~2;{^l1w&QZHR_9;3)i@QRiS6V~KL_hVD zbF^Xt$+Vy3i6G>!C8CoEOGjtdb_@@?nDwvYYRa-l+wmNpNwJXDDj^^suvc%cIWoG1 zHJeRF-D&1%G)HJtK$zZXUZ;7cXA|6&9(n29(FQonF0}IbSa%WQH`U}j_lMxKAq);3!3NF4#85-bt(?P>H&6B@*0Q ziP+)->@z#8m>s>-F@J1$9!y?X340D5R(wnd{t3=ZDr|&2k&UKR`SR=2(-$V^vvpiZ zOBE*xnf}0~gcBR9h?#9Vl_Hbxipt7I_7;|v`*D7Suj29wp^LSxA!a8JzSrF2@5cR} z_)V>oEW$(g#!74{j9{?_h8I12ZN!@nnS^Y;ADDj*L12hBGSukGrpTTk6m0(p1yy{` zG=ehYM`0-E3$1NWCuu9(Ms1{^1E>TfN;jWc02`QNDz1-X?`deG80XQk>wK^FhiH$J z=d}gHqIta0(T5krS0Y-S>9@s%%;J&=sOyx&C{DNyMX2ML5rTxDN8J_lGGa2vpn)h? zHd34C=8V;Z^)EHDHHjdM5{WY-Rx_eC2T@_wG`f~KMxz!R3K#)2Z%K6&=OaGVwu+-R zR^yLY6%T7A4)BOPdLD$CV=nP4cg`FR$*u2(=k7s~Wac2`;pijvt_Uz`k`nKUx?kG? za`D$CSTm~dRH>Sgj%Xb)TLZU<7-Z~SNJjU%P|a}2VF_QOhnzJ|Ls4_y7!q_pgcJHD zs+}%}zyW8PvdB1OEf*xTHin1iK|Apj?ku9+fijCVt7-HoBl{&Ha0H!k;>kHp@=9bn z4vKq)yI-rNQ!~=)3NNdfV9v5Ffd&=K^F`QurllS-I=K{7CeWZk_>kI~&80yd5mu+` z!o`9LxqG%p)VCWn_#L}nn06*V%4+=)jNvs{P##-}7b*&P>eR}L&idX0h*=NV2>-GE zmAW4IYFJWh4chaQv;RFclrV%bubOEg_&aiwKlE~M-(x9{;mtZ2P;%IWBR3Fz4AF}0 zUENu(`lINh3>ZVAiXe7B21gxG1 zBXPbcGdjDS+EBphq_!+`#?SF=V>~6W2@Wsde3h4EFjP8c1O?|;LPvU!p4ZsV^L5L*It4KtH(}o=;(^K)b#m4HpXz_YT{*l~9P7|MjXes_a+>(`oM%FgeWcRg!SUQ$q939xyW&x{t5ySq5DI}Y*x|9YXHL$w zcn!a9(CZJyQiy_xo4cL_AqWS}FW%fxdvDC${UnI$f#RHzt~~n-DOuHH-AL%5sfCl< zqUK=-!rwXeoaPL6qNY*R z@%C7Rg{~!ZjdHM0=TI>IC5!#d;9H#GW+mA|9*}=Unzym;(1Q}x0lL(-$<+E|2q-S*K;}4Ro#I~g~;KToZNZazg#diHIG3vitiaGoMN;f-kse^=Ra zl6UK!`>~cI+y5>DJXOr|<`-~j1=kdoF`Gi>Q=;e?*k#NXe(dK~%8<1hDsWJ2Ia_hO zKbRlgAfR(Hf@K@AI*lW{&!}zF8KR=?#O%kz4-i>>k9KyztBB&l;I+MBssEh9y=N^*NVcJ@@hP`7%&s6e=^p zL$4h#S)`#6Iww8a*UVU~66+wRIkDyOHB>Spn(|8#XiDf>!szxeDUZIK)s;k6+Xp5u5Mn4vCB6G8!0t>1 zteusYe`?@w$%9+B8Cnd1D#OYZNtn}NK&x`bQFnd=GNw%%@!D9Z{{WyKL7 zK$U79NxM$mmrw{}9}hxFUgqbz58@Sc>30!jhfbDrpXoMqybg4==NT&^hrYMhI-XU$ z0`E&AMAsA9A>!1$sQf8$J6+=YWB#o~Zxm7V!cIop#P&A7)%JFk-iDBT%;q&5Oq>o?4;cSvl7`qhkJ{ecC-y7^e1{2$ucBFJ<`)M8yy zU6G1$E1F^0!JkeOa%3oFkC&?sbcW+gAc1$u!iT|uaI6=n2~Bsx?5yMm0Pq&klEKmD zlwr zL3XqRd&YVL8#%)-vATrBOyU{{Wad1uHS8C|9NK7ti`#u=sx(LTR+QRlLQ{uR&x@aMYzOnOgcy<4r?Vesq% z56;+fQo1r!R@#ebzF~!VR+`yP4`roF>L_ebhV#>y(La+;EqA-l9s7&oc*h;o&h^!> znF*n&E4;7}?#w#p`~5;nJI4>m52%Tt^{Uf{=2~VloTiIJ&ddx-!VI*7xGkVxS!@5= z-hF@85gakYrjud?o@heXn)0?;_?uhi+wbr{tbTUwfJ1<3PMqCY2_CrQ}DC z90*;xyG*JMDUDWEn@aO#X!-88?RLmwe9NzSK7_WdgjbO+M4g21PO0{j91|-jLHm~{ znM|(ehjKo1&oyF0N{cgFn-`ulz2spF7xV@ly!}IMvPaR?b!Sdr+=|JGuAJFREPn;G z_42V!*O{Ei)a<)$C$@->PXJ-kkY&gG8B1zIUB>q~n?q{vq~}a5(so~pHv#!A+7(;9 zu>zfbA$amQd`KgK*-G|Epbt)mVFLD~Y(cF!d{TQ%HXuz22NyWY#zdib1 z2FhnWR-j^^+$Jl^r|n#%4(FvipF7#Y5_3T3PQ9_3XptDdXX!T+8c}+e+go$q0po?4 ze_VaBz6jT$=}MD3kaeC|*Ice)bXp@hB&MgYZ{Zgh|JN2L#0yfp!?QduqUjohSYR?=a29Q@!ZorkK+tF8IZZN6ol^%ap)7CHEaDVEx-Z)rDT zyWR3?yHk{nO~%>tXX1Z6ah7&!kfDCD3jy5;5uMC>bgFV&aJQDzf!zG$*Dj#zcV@ls z>4Gt4;0HDi;Ko7~f8p=drHP+3>>93l-NtIIex_XO+I1+15&YHBsREZ|i(F71y|sMf zM?|zK-7EC9Luj}(yS`SBRA}jqk}asJW6BRB5-h{d6YG=mn8I3TSo&reFZGe>SbP0E zM>yzAUNtRO;q-KTr&r{jvcuj)s|xCVAI;khrUJXmypcx3^ou1B)JwC|r-lX0Umg=C-oBq}{I$vP`hp!z#3{qr-C^-S_11-s=U$wp z3#A;>Po&W6{0J5e{PSy%XHt1_+zrK&KQOZr<$SmJ9`{NWH`P%N7TaXd(+k@;0oBKU z_FT*LzuM#AYX=klS+%dJ z^YuW|;~eJB10CTGPELc?n?|}vlBc=7HKZ3(^Vf+7z3VUU4J(cO{w78AfLYfk zZ?s-=z3Gg46;QjunC~Dmvaa`8zm1QEVtiTt%8^3&n=*NsYGca4r$$WH87*mr<3gWp z7KhjJBr4j3PkGkn_t%RLxnV05{!nC8l$~m6u-^Mow)CWmCEk$vwW0mQCnImmpNROQmjY%|GEOa~SdIY4?%dG&cmKcNj7cxnVBi?o>C0FA?4} zYk)Wlrw(GpJ85C(!umt|O3EjRPvq(*m<1_Fz0bCfP*q|FFqZ1$&U*VYOZ>`_3qK=Y ziif;4!dpZ&BK0QtEIAdD`)V5{lWw#9rk+C+?H2ZB41M;~yYw9Xt?blEPmRz!PKrvm zR5C5LnD9C9Xdw56w6OjRo@AxlS6p+u`;+^v%Y3Q^AM;$s#M0dJ<+P;Hkr%YI?#d=( zJHH)So9@Ij5@lX(Q`ywxn1KO1i?n4Y6y2$YUw_o*yXgSM^UAZkv*iMpcYoiL-+0Oc zedqR>yQh^!f?Kv#7HxF|Rg=CN_2S;2W0cAUoROt5u6yloR%F-9XDLsvD0!dhu)7D| zXYbfImlxkn!pwmZa{kD^#eRB7w(?zFHGB?o&f=cQ$dV)9g(?lb_hg5fU6DFTuo23= zikmui^vX^*qVV+S_%`;7h2e!8Zx?==E)6s;^h?WE&_t7+&Bn{lp3i*ld{IH8>X}*i zgM$&WIT_&!#}h!e$rZuJcXeU;Rnk=s(05l4BCqu3)RGJ<6(09i=9~%&jC^6}q#-vr zEb6}iN$r%Z5EQxIr_d1Qb41|n+qaNSDwP*O4<_I6gQ>M3Vu&??ro zztiSn%_D-h9~hP;(b#8(@=@N*3OGIHs6!!opM)dbh%34r6-|mP#XvTUul+zA@5ofBkCP!kGif?#sb4FcEOe|<+&eG)xA`$ZBDCp)@pt7yldBM#cd3R6sj;>|KHw_tahq_0} z+cw_|`1#b;B26l1EN$jR-ln$!009M0-7O9Q7CSFIZQVo+M_eyr1rN7?$&Lo*)WG?# zQq3GnEBzRtezsSgFl1ds%5d^aP6FkmnZwqESzqWRowE3%e0|$g^_S63{n?eEe>Vz<(>v~QBNmuK_1tGz)_q^j>etz=zBN>oI`?AjW5dr;ze-{(~zFi=GtUh+C0I}^~MV+tf zI|}mQLLIc|Kb^yW_w=do|7qEMKdgxe#Dcqx*;FS6D!+b{_vPoXq!z^W3GYgU?RrgvN_sCp*Z_ZYwoRN`{l|!4} z*71lq_CD^Ew5)u^rw51ZM@o(#s%TS?`}9D*^h40n%7c$mis{Cmu0|wh-S#LsX4b_r zwFxpA9cz53n$p;a_AjkY5#P)5lYK2)cS49f-akWhmlc1rEr+o}cdOw{k671z*Vn&# zTr{)B^eT85T=LZCldX`|&0UXf*Zmwys2e)fw21-%fJ@B1n4b}EG_D@9cS9<~mACP< zjRDpOW>|m^oNrb0PI` zS59W?nxP6=+r<4{i&YcqfSPY4kb$6E` zt|-{5N+l3gPt04_sVj8vYI^yy+SQfM(bw4h_zA_V3o_XUw}%$DcwKEwfP6ii(Io?{ zcyPwOTX=30ZSc}-5OKUHZ0LINqkd~kORu?$T{e#@1j^1ATWOv!dG-8(jq%}PD{O-G zh0qB{t+T0*Y~ki~q`oIZDe{EpaIQ_&he@826%^jCCRMvyar#8iDaSN&N*+rKqama%qn*tu?uy?{|-S?z-ktGITE7n7+N_jBlSs6u=NSHuf@all^mJ zm6zJ~?Kz)8`~zE?p()jrjEIWQ`d9TS3EG4UD=s%}gTD4zfE|4tmg93S<_^*_zTH1+ zPfg3IF}x%^ZQ3Fo|0(T(9whl1pRVuKvruC`8(}hJRF740eB;Tktg1wOSsX5*JNa?m zKR!f%u%5^vsY%fYY1sq5M>G}XE1sbl{b8>xzFb{8%ff0Ai;u#FcC=>7R6UhdlXvw5 z>C2>CD}@~g4C~F#~JX{nrrQFMcYZI0rvOJ+4=wg@_dx8 zB^WSN>r*Ud@yrB-~!#vP7u zIt!z&i&xrLh~!_CYiN=#1GPe0*(zufrd$~yM*4>)C@70llJ@%Y(0cjf(3F)Ax3JK_E%LvSpX1cq?mSAX-{lDjFIU z#2e*cjgYzLd1BWpv)}()iE1D>iOras&%4~*z+gKbwtxjiUMn0Zw~$D&$lKIto6Gz>UMy`A{fq2!^>&nV^aYXoMW3vaJx#5rfgsgDDDB{Dq!{_Q0X9U~9^{X05V zp)ppbR&O0#Z~phsdz?H(2`0xrE|A{*MF6)S{x1}iC#Feh<)+=wqgh=sGih`l2nHB- zEiP972e#X8Q1M66b5jpAWTO16{iMM|w=0*Ar>3O9ubZG9vKnVo_wG0D1HYQDPODkF zeehsyR8{q0tvdv0y*|-0bjVCHgZ#-oex9|^IJ^~oz~OotxuYz5huNkn=A}WmAK&$4 zw~3eX@lWZYo;wX>;b>#!;|Ue7oi7592X>j1SOr>v+-AQy$UL$R(mt3zZq?Z6JNdM@ zhP<3DQ{`r)*Pgw*`bqT{JuUSRRz#Mu(PBXbFdFDU(r%`w$+T^l%=;GUMfT`z=SzyuH2ut2NxA8U^w(JMc-^0iJMN77o9ntb9gU2Hlr&7I?+5V;zh8LjbZ~rOb+wk>OvUlJ zO}aq zne>N`+lKa$ibtMxCx5o(KK$=g&Z+Zd|HmrFE6E};#Yd#RkxBncJ zUpa43jem$bBn!+u0fyxDV9xE3dT3qeP;upxaYQ`&TBk~=-rz&2h8O$d?@UMODO#$5 zzYp0tr^qUFtI10Mj{HZj<$p)zl)pALi}xX)J8aWuW^z|kPU`+Q-ymV4MKBkIH%bK**yRAa&CVJ#_-s{Iyx8mS7# zx0IzOpM%`2J0^YO<)D?j)W$(JiTW$@G0CqR8%sC8=eJ}%AEIk{ATRk*;faB`SF{f) zV@*Cn(IUGa<{ToVz>vXaw;<2J-;kyb?lu)SS;n;Gp6P2XYfV1x#&tfz6Z;OjGi7bl z$ab-X-CA|RHA0k4suBR`_wEomd4_kl|7l~|N2ZMFn%goUt;5Sq3gHzb9kSH*_D5pEZl2BY)uP}^1XomT$Bo<3 z|M;9ML;H3Z>}?J(&U*oIUNBmw**(W-ulqTy?5h1|XT-J}g#H1h)E#p$egbp0ai|+K z337eZ!P#=8dUVv`KcTLgg>lQlLS6ReGycZi!2e0x&+@rPa8maz*=nM}EdkM8z+l~OuhN`!&%sz5Y0nz6G!O93 zGPR4$eA(ydjU%0wg12T_k;*}vDH;&>zJEIp0q?T_*=o)eHhvUknAF_jzhM)rUVSb% z^uA_2l9d$~)c}u=tJ@2#e^GVO>G=E`=Ng@}vWEJ*%W9&V7b1N=)WTmas3rH1njw2U z767WoLU2E0yvP!zPl*84oTPMz7 zKsL?p4zByVMO$lxzmEXt)v*0h_^IW(sU z#n=gy@*G!;zd?LQ^f;8a$mOgSnU}B<;tfL%r52Q%Lzh#*0p?D(f>WgZOs--w8VwG_ zm-}Q%7kzk}63h%LdT33Q;vsO>ueGh`-02@G6m4kFj1I=f#YmO>9PJ!7ld2@*tOX9g zsn{02Vi|9@4Y|ASWkUDNvf1v}?4{V^ z^?Na2E_~AO{ReV?BOD({%mr@^_{sASu45oi*c{jfl~iG8FnwicrM`IU<5D-jn;