From 136f725d280f5bcb4d53627991ca04fd0c18d177 Mon Sep 17 00:00:00 2001 From: Zichuan Date: Mon, 17 Sep 2018 00:11:33 -0400 Subject: [PATCH 1/5] Finished naive --- src/main.cpp | 6 ++- src/testing_helpers.hpp | 14 +++---- stream_compaction/CMakeLists.txt | 2 +- stream_compaction/common.h | 5 +++ stream_compaction/cpu.cu | 54 ++++++++++++++++++------ stream_compaction/efficient.cu | 10 ++--- stream_compaction/naive.cu | 70 ++++++++++++++++++++++++-------- 7 files changed, 119 insertions(+), 42 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 1850161..04dad19 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -6,6 +6,9 @@ * @copyright University of Pennsylvania */ + + + #include #include #include @@ -29,6 +32,7 @@ int main(int argc, char* argv[]) { genArray(SIZE - 1, a, 50); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; + onesArray(SIZE, a); printArray(SIZE, a, true); // initialize b using StreamCompaction::CPU::scan you implement @@ -51,7 +55,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 diff --git a/src/testing_helpers.hpp b/src/testing_helpers.hpp index 46337ab..d6cc4e3 100644 --- a/src/testing_helpers.hpp +++ b/src/testing_helpers.hpp @@ -1,8 +1,8 @@ #pragma once -#include -#include -#include +#include +#include +#include #include #include @@ -69,8 +69,8 @@ void printArray(int n, int *a, bool abridged = false) { printf("]\n"); } -template -void printElapsedTime(T time, std::string note = "") -{ - std::cout << " elapsed time: " << time << "ms " << note << std::endl; +template +void printElapsedTime(T time, std::string note = "") +{ + std::cout << " elapsed time: " << time << "ms " << note << std::endl; } \ No newline at end of file diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index cdbef77..6ad13f6 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -13,5 +13,5 @@ set(SOURCE_FILES cuda_add_library(stream_compaction ${SOURCE_FILES} - OPTIONS -arch=sm_20 + OPTIONS -arch=sm_52 ) diff --git a/stream_compaction/common.h b/stream_compaction/common.h index 99a1b04..bcca058 100644 --- a/stream_compaction/common.h +++ b/stream_compaction/common.h @@ -13,6 +13,11 @@ #define FILENAME (strrchr(__FILE__, '/') ? strrchr(__FILE__, '/') + 1 : __FILE__) #define checkCUDAError(msg) checkCUDAErrorFn(msg, FILENAME, __LINE__) + +#define BLOCK_SIZE 1024 + +#define PLUS_OP_IDENTITY 0 + /** * Check for CUDA errors; print and exit if there was a problem. */ diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 05ce667..99acff7 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -1,15 +1,15 @@ #include #include "cpu.h" -#include "common.h" +#include "common.h" namespace StreamCompaction { namespace CPU { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } /** @@ -19,7 +19,10 @@ namespace StreamCompaction { */ void scan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + odata[0] = 0; + for (int i = 1; i < n; ++i) { + odata[i] = odata[i - 1] + idata[i - 1]; + } timer().endCpuTimer(); } @@ -30,9 +33,15 @@ namespace StreamCompaction { */ int compactWithoutScan(int n, int *odata, const int *idata) { timer().startCpuTimer(); - // TODO + int non_zero_idx = 0; + for (int i = 0; i < n; ++i) { + if (idata[i] != 0) { + odata[non_zero_idx] = idata[i]; + ++non_zero_idx; + } + } timer().endCpuTimer(); - return -1; + return non_zero_idx; } /** @@ -41,10 +50,31 @@ namespace StreamCompaction { * @returns the number of elements remaining after compaction. */ int compactWithScan(int n, int *odata, const int *idata) { - timer().startCpuTimer(); - // TODO + + // allocate a temporary 0/1 accumulating array + int* temp = new int[n]; + + //for (int i = 0; i < n; ++i) { + // temp[i] = 0; + //} + temp[0] = 0; + timer().startCpuTimer(); + // scan to 0/1 accumulating array + for (int i = 1; i < n; ++i) { + temp[i] = temp[i - 1] + (idata[i - 1] != 0); + } + + // use temp to map to output + int count = 0; + for (int i = 0; i < n; ++i) { + if (idata[i] != 0) { + ++count; + odata[temp[i]] = idata[i]; + } + } timer().endCpuTimer(); - return -1; + delete[] temp; + return count; } } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 36c5ef2..fd1d622 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -5,11 +5,11 @@ namespace StreamCompaction { namespace Efficient { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } /** diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 9218f8e..27e43aa 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -1,25 +1,63 @@ #include +#include #include #include "common.h" #include "naive.h" namespace StreamCompaction { - namespace Naive { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; - } - // TODO: __global__ + namespace Naive { + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; + } + + __global__ void kernNaiveGPUScan(const int n, const int offset, const int* d_data_in, int* d_data_out) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx >= n) return; + if (idx - offset < 0) { + d_data_out[idx] = d_data_in[idx]; + } + else { + d_data_out[idx] = d_data_in[idx] + d_data_in[idx - offset]; + } + } + + __global__ void kernIncToExc(const int n, const int* d_data_in, int* d_data_out) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx == 0) { + d_data_out[0] = PLUS_OP_IDENTITY; + } + else if (idx < n) { + d_data_out[idx] = d_data_in[idx - 1]; + } + } + /** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ + void scan(int n, int* odata, const int* idata) { + dim3 blockSize(BLOCK_SIZE); + dim3 gridSize((n + BLOCK_SIZE - 1) / BLOCK_SIZE); - /** - * Performs prefix-sum (aka scan) on idata, storing the result into odata. - */ - void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); - } + // allocate memory + int* d_data_in; + int* d_data_out; + cudaMalloc((void**)&d_data_in, n * sizeof(int)); + cudaMalloc((void**)&d_data_out, n * sizeof(int)); + cudaMemcpy(d_data_in, idata, n * sizeof(int), cudaMemcpyHostToDevice); + timer().startGpuTimer(); + for (int i = 0; i < ilog2ceil(n); ++i) { + kernNaiveGPUScan << > > (n, 1 << i, d_data_in, d_data_out); + std::swap(d_data_in, d_data_out); + } + timer().endGpuTimer(); + // for readbility + std::swap(d_data_in, d_data_out); + cudaMemcpy(odata + 1, d_data_out, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); + odata[0] = 0; + cudaFree(d_data_out); + cudaFree(d_data_in); } + } } From cc061ea1df32bddb0d98a279211d43bd00f7e3f0 Mon Sep 17 00:00:00 2001 From: Zichuan Date: Mon, 17 Sep 2018 00:38:24 -0400 Subject: [PATCH 2/5] finish naive --- src/main.cpp | 4 ++-- stream_compaction/efficient.cu | 32 +++++++++++++++++++++++++++++--- stream_compaction/naive.cu | 2 +- 3 files changed, 32 insertions(+), 6 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 04dad19..8b95447 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -68,14 +68,14 @@ 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); 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); diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index fd1d622..ca9e38c 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -12,13 +12,39 @@ namespace StreamCompaction { return timer; } + + __global__ void kernEfficientGpuScan(int n, int *odata, const int *idata) { + + } + + __global__ void kernBuildTree(const int n, const int* d_data_in, int* d_data_out) { + + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); + dim3 blockSize(BLOCK_SIZE); + dim3 gridSize((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + + // allocate memory + int* d_data_in; + int* d_data_out; + cudaMalloc((void**)&d_data_in, n * sizeof(int)); + cudaMalloc((void**)&d_data_out, n * sizeof(int)); + cudaMemcpy(d_data_in, idata, n * sizeof(int), cudaMemcpyHostToDevice); + timer().startGpuTimer(); + + + + timer().endGpuTimer(); + // for readbility + std::swap(d_data_in, d_data_out); + odata[0] = 0; + cudaMemcpy(odata + 1, d_data_out, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); + cudaFree(d_data_out); + cudaFree(d_data_in); } /** diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 27e43aa..153e5bc 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -16,7 +16,7 @@ namespace StreamCompaction { __global__ void kernNaiveGPUScan(const int n, const int offset, const int* d_data_in, int* d_data_out) { int idx = threadIdx.x + blockDim.x * blockIdx.x; if (idx >= n) return; - if (idx - offset < 0) { + if (idx < offset) { d_data_out[idx] = d_data_in[idx]; } else { From 3fc681a1c203c540c56f3503bc3bf2b4e70ff69c Mon Sep 17 00:00:00 2001 From: Zichuan Date: Tue, 18 Sep 2018 22:52:03 -0400 Subject: [PATCH 3/5] Finished all code --- stream_compaction/efficient.cu | 240 ++++++++++++++++++++++++++------- stream_compaction/thrust.cu | 38 +++--- 2 files changed, 209 insertions(+), 69 deletions(-) diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index ca9e38c..e217897 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -2,65 +2,201 @@ #include #include "common.h" #include "efficient.h" +#include +#include namespace StreamCompaction { - namespace Efficient { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; - } + namespace Efficient { + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; + } - __global__ void kernEfficientGpuScan(int n, int *odata, const int *idata) { - - } + __global__ void kernEfficientGpuScan(int n, int *odata, const int *idata) { - __global__ void kernBuildTree(const int n, const int* d_data_in, int* d_data_out) { - - } + } - /** - * Performs prefix-sum (aka scan) on idata, storing the result into odata. - */ - void scan(int n, int *odata, const int *idata) { - dim3 blockSize(BLOCK_SIZE); - dim3 gridSize((n + BLOCK_SIZE - 1) / BLOCK_SIZE); - - // allocate memory - int* d_data_in; - int* d_data_out; - cudaMalloc((void**)&d_data_in, n * sizeof(int)); - cudaMalloc((void**)&d_data_out, n * sizeof(int)); - cudaMemcpy(d_data_in, idata, n * sizeof(int), cudaMemcpyHostToDevice); - timer().startGpuTimer(); - - - - timer().endGpuTimer(); - // for readbility - std::swap(d_data_in, d_data_out); - odata[0] = 0; - cudaMemcpy(odata + 1, d_data_out, (n - 1) * sizeof(int), cudaMemcpyDeviceToHost); - cudaFree(d_data_out); - cudaFree(d_data_in); - } + __global__ void kernBuildTree(int n, const int* d_data_in, int* d_data_out) { + + } + + __global__ void kernEfficientScanUp(int n, int bitShift, int* d_data_in) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx >= n) { + return; + } + int pow1 = 1 << bitShift; + int pow2 = 1 << (bitShift + 1); + d_data_in[idx * pow2 + pow2 - 1] += d_data_in[idx * pow2 + pow1 - 1]; + } + + __global__ void kernEfficientScanDown(int n, int bitShift, int* d_data_in) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx >= n) { + return; + } + int pow1 = 1 << bitShift; + int pow2 = 1 << (bitShift + 1); + int pos1 = idx * pow2 + pow1 - 1; + int pos2 = idx * pow2 + pow2 - 1; + int temp = d_data_in[pos1]; + d_data_in[pos1] = d_data_in[pos2]; + d_data_in[pos2] += temp; + } + + /** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ + void scan(int n, int *odata, const int *idata) { + // make length to 2^n + int level = ilog2ceil(n); + int trueN = 1 << level; + + + + std::unique_ptrtrueIData{ new int[trueN] }; + + // pad 0 to the end + for (int i = 0; i < n; ++i) { + trueIData[i] = idata[i]; + } + for (int i = n; i < trueN; ++i) { + trueIData[i] = 0; + } + + // allocate memory + int* d_data_in; + cudaMalloc((void**)&d_data_in, trueN * sizeof(int)); + cudaMemcpy(d_data_in, trueIData.get(), trueN * sizeof(int), cudaMemcpyHostToDevice); + timer().startGpuTimer(); + + dim3 blockSize(BLOCK_SIZE); + dim3 gridSize; + int pow2; + // go up + for (int i = 0; i < level; ++i) + { + pow2 = 1 << (i + 1); + gridSize = ((trueN / pow2 + BLOCK_SIZE - 1) / BLOCK_SIZE); + kernEfficientScanUp << > > (trueN / pow2, i, d_data_in); + } + cudaMemset(d_data_in + trueN - 1, 0, sizeof(int)); + + // go down + for (int i = level - 1; i > -1; --i) + { + pow2 = 1 << (i + 1); + gridSize = ((trueN / pow2 + BLOCK_SIZE - 1) / BLOCK_SIZE); + kernEfficientScanDown << > > (trueN / pow2, i, d_data_in); + } + + timer().endGpuTimer(); + // only need copy n, no need to copy trueN + cudaMemcpy(odata, d_data_in, n * sizeof(int), cudaMemcpyDeviceToHost); + cudaFree(d_data_in); + } + + __global__ void kernValueMapToOne(int n, int* d_ones_out, int* d_data_in) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx >= n) { + return; + } + if (d_data_in[idx] != 0) { + d_ones_out[idx] = 1; + } + } + + __global__ void kernCompact(int n, int* d_indices, int* d_data_in, int* d_data_to_compact) { + int idx = threadIdx.x + blockDim.x * blockIdx.x; + if (idx >= n) { + return; + } + int data = d_data_in[idx]; + if (data != 0) { + d_data_to_compact[d_indices[idx]] = data; + } + } + + /** + * Performs stream compaction on idata, storing the result into odata. + * All zeroes are discarded. + * + * @param n The number of elements in idata. + * @param odata The array into which to store elements. + * @param idata The array of elements to compact. + * @returns The number of elements remaining after compaction. + */ + int compact(int n, int *odata, const int *idata) { + // ugly implementation of reusing scan code + // used for mapping to 1 and as the scan result + int* indexArray = (int*)malloc(n * sizeof(int)); + int* d_data_in; + int* d_compacted_data; + cudaMalloc((void**)&d_data_in, n * sizeof(int)); + cudaMalloc((void**)&d_compacted_data, n * sizeof(int)); + + cudaMemset(d_compacted_data, 0, n * sizeof(int)); + + cudaMemcpy(d_data_in, idata, n * sizeof(int), cudaMemcpyHostToDevice); + int* d_ones_scan_result; + + // calc true values + // make length to 2^n + int level = ilog2ceil(n); + int trueN = 1 << level; + + // set scan_result zero + cudaMalloc((void**)&d_ones_scan_result, trueN * sizeof(int)); + cudaMemset(d_ones_scan_result, 0, trueN * sizeof(int)); + + // useful constants + dim3 blockSize(BLOCK_SIZE); + dim3 gridSize((n + BLOCK_SIZE - 1) / BLOCK_SIZE); + int pow2; + + timer().startGpuTimer(); + // do ones + kernValueMapToOne << > > (n, d_ones_scan_result, d_data_in); + + // TODO(zichuanyu) make this a in the future + // scan + // go up + for (int i = 0; i < level; ++i) + { + pow2 = 1 << (i + 1); + gridSize = ((trueN / pow2 + BLOCK_SIZE - 1) / BLOCK_SIZE); + kernEfficientScanUp << > > (trueN / pow2, i, d_ones_scan_result); + } + cudaMemset(d_ones_scan_result + trueN - 1, 0, sizeof(int)); + // go down + for (int i = level - 1; i > -1; --i) + { + pow2 = 1 << (i + 1); + gridSize = ((trueN / pow2 + BLOCK_SIZE - 1) / BLOCK_SIZE); + kernEfficientScanDown << > > (trueN / pow2, i, d_ones_scan_result); + } - /** - * Performs stream compaction on idata, storing the result into odata. - * All zeroes are discarded. - * - * @param n The number of elements in idata. - * @param odata The array into which to store elements. - * @param idata The array of elements to compact. - * @returns The number of elements remaining after compaction. - */ - int compact(int n, int *odata, const int *idata) { - timer().startGpuTimer(); - // TODO - timer().endGpuTimer(); - return -1; + // compact, only use useful part of the array + gridSize = (n + BLOCK_SIZE - 1) / BLOCK_SIZE; + kernCompact << > > (n, d_ones_scan_result, d_data_in, d_compacted_data); + timer().endGpuTimer(); + cudaMemcpy(odata, d_compacted_data, n * sizeof(int), cudaMemcpyDeviceToHost); + // count how many nums + // cpu or gpu + int num = 0; + for (int i = 0; i < n; ++i) { + if (odata[i] == 0) { + break; } + ++num; + } + cudaFree(d_compacted_data); + cudaFree(d_ones_scan_result); + cudaFree(d_data_in); + return num; } + } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 36b732d..9693dca 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -7,22 +7,26 @@ #include "thrust.h" namespace StreamCompaction { - namespace Thrust { - using StreamCompaction::Common::PerformanceTimer; - PerformanceTimer& timer() - { - static PerformanceTimer timer; - return timer; - } - /** - * 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()); - timer().endGpuTimer(); - } + namespace Thrust { + using StreamCompaction::Common::PerformanceTimer; + PerformanceTimer& timer() + { + static PerformanceTimer timer; + return timer; } + /** + * Performs prefix-sum (aka scan) on idata, storing the result into odata. + */ + void scan(int n, int *odata, const int *idata) { + thrust::device_vector d_data_in(idata, idata + n); + thrust::device_vector d_data_out(odata, odata + n); + timer().startGpuTimer(); + // TODO use `thrust::exclusive_scan` + // example: for device_vectors dv_in and dv_out: + // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + thrust::exclusive_scan(d_data_in.begin(), d_data_in.end(), d_data_out.begin()); + timer().endGpuTimer(); + thrust::copy(d_data_out.begin(), d_data_out.end(), odata); + } + } } From ff1e59fa829aba5c588b85c17900e2172b73be20 Mon Sep 17 00:00:00 2001 From: Zichuan Date: Tue, 18 Sep 2018 23:45:25 -0400 Subject: [PATCH 4/5] readme ver 1.0 --- README.md | 99 ++++++++++++++++++++++++++++++++++++++++++--- img/block_size.png | Bin 0 -> 15370 bytes img/compaction.png | Bin 0 -> 17695 bytes img/scan.png | Bin 0 -> 19180 bytes src/main.cpp | 2 +- 5 files changed, 94 insertions(+), 7 deletions(-) create mode 100644 img/block_size.png create mode 100644 img/compaction.png create mode 100644 img/scan.png diff --git a/README.md b/README.md index 0e38ddb..9b61c8b 100644 --- a/README.md +++ b/README.md @@ -3,12 +3,99 @@ 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) +* Zichuan Yu + * [LinkedIn](https://www.linkedin.com/in/zichuan-yu/), [Behance](https://www.behance.net/zainyu717ebcc) +* Tested on: Windows 10.0.17134 Build 17134, i7-4710 @ 2.50GHz 16GB, GTX 980m 4096MB GDDR5 + +## Features + +- CPU Scan +- CPU Stream Compaction +- Naive GPU Scan +- Work-Efficient GPU Scan +- Work-Efficient GPU Stream Compaction +- Thrust Implementation + +## Performance Analysis + +### Block size analysis + +![block_size](img/block_size.png) + +### Array Size Analysis on Scan + +![scan](img/scan.png) + +### Array Size Analysis on Compaction + +![compaction](img/compaction.png) + +## Output + +Array size 2^26, block size 1024 + +```shell + +**************** +** SCAN TESTS ** +**************** + [ 1 1 1 1 1 1 1 1 1 1 1 1 1 ... 1 1 ] +==== cpu scan, power-of-two ==== + elapsed time: 1535.85ms (std::chrono Measured) + [ 0 1 2 3 4 5 6 7 8 9 10 11 12 ... 268435454 268435455 ] +==== cpu scan, non-power-of-two ==== + elapsed time: 594.798ms (std::chrono Measured) + [ 0 1 2 3 4 5 6 7 8 9 10 11 12 ... 268435451 268435452 ] + passed +==== naive scan, power-of-two ==== + elapsed time: 510.046ms (CUDA Measured) + [ 0 1 2 3 4 5 6 7 8 9 10 11 12 ... 268435454 268435455 ] + passed +==== naive scan, non-power-of-two ==== + elapsed time: 510.037ms (CUDA Measured) + [ 0 1 2 3 4 5 6 7 8 9 10 11 12 ... 0 0 ] + passed +==== work-efficient scan, power-of-two ==== + elapsed time: 175.304ms (CUDA Measured) + [ 0 1 2 3 4 5 6 7 8 9 10 11 12 ... 268435454 268435455 ] + passed +==== work-efficient scan, non-power-of-two ==== + elapsed time: 175.151ms (CUDA Measured) + passed +==== thrust scan, power-of-two ==== + elapsed time: 28.8416ms (CUDA Measured) + passed +==== thrust scan, non-power-of-two ==== + elapsed time: 28.8394ms (CUDA Measured) + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 0 1 0 0 1 0 3 1 3 3 0 3 1 ... 3 0 ] +==== cpu compact without scan, power-of-two ==== + elapsed time: 708.621ms (std::chrono Measured) + [ 1 1 3 1 3 3 3 1 1 1 1 1 1 ... 1 3 ] + passed +==== cpu compact without scan, non-power-of-two ==== + elapsed time: 680.761ms (std::chrono Measured) + [ 1 1 3 1 3 3 3 1 1 1 1 1 1 ... 1 1 ] + passed +==== cpu compact with scan ==== + elapsed time: 1471.92ms (std::chrono Measured) + [ 1 1 3 1 3 3 3 1 1 1 1 1 1 ... 1 3 ] + passed +==== work-efficient compact, power-of-two ==== + elapsed time: 213.044ms (CUDA Measured) + passed +==== work-efficient compact, non-power-of-two ==== + elapsed time: 212.931ms (CUDA Measured) + passed +Press any key to continue . . . +``` + + + -### (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/block_size.png b/img/block_size.png new file mode 100644 index 0000000000000000000000000000000000000000..afdd5f5d005bc81f77452ddde7239253883632ea GIT binary patch literal 15370 zcmb_@1yqz>zwf9BC@Bg^cS$25E!`pAC=59ah?KMf!mE^Y2}n0c3AarYnz7T=wM*MM|L}yE<;=7xpVY1Q{S?*l>|g6czL|AS9y&AI z5rtExo{vJbhA6vL*;k8;iz)dlyE}`Ch(9V3K7$geVLipXVWp$<07EYKft9M2CUeql zXtZbLjUDs!?zk@TOj?KOIgMhT_)x@{*HNWlpVQTtkvA$334uH*r|G{9fdml#Z$AW8 z)zk*=P~t-{GVVLOxy>M83CT^b_cB96RdjXZ{N4=?4*vS}Tu4X=E&#)!8Y;JJTT+^- zbL*PCOUh;Pk+{^Xq3#WGV{`MyYxjK;wzQnrnyB}zJUF^f7PJy_3nFE!;^9$KR#p}!_2R{g zwY4=q>eKV{F8|~085qw))qt0MY2tZ}E(gB{wx()|-c2UuE8sMDbSNpA8%i}l5f&C^ zVv=jMJ$2I5>_0jAx`>rg8x(|L=dYp?I`a-=J^+gj60l3x-QBHt)DCey-M4~g39MQd_o&8q~sim zb4NOW@S7H9MwkWP8Kv;6pU~Rj07pm1w6wJ7Siu0zgQKIu%$E{UuXle*o!aT^C-`MR zGeSZ_z|6!)L?c&wQ>SWO@`66?bf);#buGF>cxipzJ`@L$XJ}yX)%)1- zFk4Go`%pt!z)aa-4tYUFP8X6j3mFDJV2S#mQ-v%*fZf-qYiQgP$!e1vaGQxwA-Pi8W z=akj`k)6tSmqyahq{3V|lhe2w7Z-PYi^(%$|EcF!#_t_5ELoYE&^NPHb}=oqu{Vr5 zVFCTw2O*oFIA>{eAzStFv?Ek`<;=otY_3 zPP;YN=!@*OjdI|{wrV1BU{U|{(UOHCD&H0ct9%$V_{dH3XnQ6lE{=&HLA>(^!lBX^ z<4}r~ife)ySve{wJRJ>v^>g`USFTQ*Jqy7BUQBA5_s7DnWlj5@>X{@++AeJ@ol+dW z`&bo8&{&yy!g4eo>W(y-u~mQUTW$u1EF^NO_5BmyTT7<(98tXbD>(;y$Hz{Nj-!5$ zMlu?i3SSJWZ!v$qxbQ(MnNmOxVg;C}B6bG4rGZyh*=XEF*<_wuKy&F_e%W6}*p zSW8Z3b*FDq2KF!cAiTckm-2XyzT(JEV6;2fD6&vV)kOGWAZTOPwXG{EE8$g~yu8n* ziuD*qM@I`A5hOu_mLd35(KH3kUJaRG0;Wl5L!Ha2I;_X9N+C&dLIG*WEFhpuA}A=Boc=b}Y_h?-rb9ma%8IJs zVv|}5A?kUH@oEw;BRCk-t_c)}*tPxSfsmqbTAg(6Sn}fxk-uj;Dk^GUbZ%-2cA53p zU)U>mk}pUeA^8B$w6RxP*k1X+55F+F=5X{cSVf0D5t8^fp>UWL@K1a?F;><>|@E zCvk7RWH0oYn?QE2BsNnG6j9cAX8%1euRYT0{1=bwT3Gb2xSYeEbct+SAU`Ki_(el^If3@kOHe@|a zKdscGk>nAU`CNLhF!JfRM8FwRzN@RtNL0T~#Un<~Jf?K6x2my0*b3@0i2u(K} z8Ar!biOA(1B<|?t<)wYb2LrT0ZgO^Gk?KrBS*pgy#)kHsSC(AGjKf2)i`|clClNY* z%*!hmftql`Sr^+>Zecn%7j5mBa9W=rg(R#D9a9%J_Ql0TM8ojt=;T+JPQzGK`#^R7 z=xA;eZ20IN|0!k0*~37D#zxn7Hy4*7{{y+nf*A?R!7s_l2JfGz^{*u2` zQ&tvnUI@Zf-Tw%R{P4(#-4AU8gTy*QCM#KK>42ZvI&|yN8i3H8#kvB7+I)7zb*tf2 ze|EeblWcCeiKId6ZeYoQEY8@>72bdX5gdt#braZB5s@6ciM!i;`8P=)-BGVrVJJ z$oBU4qeBSr@sn%R+#hih7|x&UEmm;56&Q<_XE-G~n4|?<=uf)LOWAk1gCZG26SQtJ zl%tZsh;Idric-hCNfuYtrzYt@INNab;^KM>jY6xUUnoY8Y@fr(A{p6o@YbiGjIW_f zd~5^EEG#)*A4TMEYT4f;t2_{rX#WO4_fcTkg`ytJ*7hT2SLtdZo)&NWu8rRZ1BwJI zob?nd#=h+Rakn`!u(zGr-i=Vv4_az{-q^Wcek5)~aL$UN-CSFnsI-X`qf+9gkYbgq zudmVM#uXN+ZZecNQ>@TLbdHVbSMwY**j^ID)_cNTlUoAm%UC@XG8^~h2X(w;kh%Fz ziTvMM?rtepZ}npQdLV^OcY{OV1{wk=<6|v<+K$!^V+cg*gs~_;|F5tx{?q$m|6{64 zsnhmN)i0H|`uVzQT)D^VczinS8{Z6Z!+X5))Gocw^999`maeM>9nL=jTIK;YT(bc@ zZ!X!*)nT^&^z3kh(g+fXB=q5QRZS0=Jn)?UdC+}x)56u&b*dUJpXHq7{><3e`1$kC zDJjP0=IPa6>g&a5X}`_PpufGng+`+>=5MB?rs5G0%+pR93J5vuEI=>}+gGii+4cI1Jb~#PoDIhXl1W zbBRS2N{HXpIL=I#TPj#v=gI~=P>sg3_XM=;QCwzL77H`8j@<8s1^UN)6~qr72<`Shn4{M21io|1h(8fq%^jMe zYnZ=WQ#(Cv!bWzd`F82|krZ;66sr@mlSI1bY)(iFCp~+<>CAzm`25g79vO6@HnLbNkGx6mx$0XW0I9x;PK)$t(00}v{y(M?F zO^P5qIz(l~E;3s6@aSlJdmBuBJlT7$iGtyRwg{@2^iE5wofYo^Cp}B&%Id%`Sy?~& z`ZS1AX-#;v>ahTx-1u@Y2&cJ!)Pmvm2y}C69?iE46@6yiL|;G14;R9=`f?W^Uk~-} z-8;ZWKJuwIb4y5=(FD65qd?EeJc2p62Fry}i9a=iMPd(Z)WAHHrF3{_vsZ zRrp_jeHoBOn=#Jn$gvX7&Ng@>279`?R31G}>T7RTR9BBLjynbaINd4G&g0l_u^=>> zDh5Q{T5(3p%xwFK0_x5M<-BDHHg`ztI9~I5!!umKZ2Fe&2Qd%5HThveuZ+`7A>^nS z#Yi_cHU@sBM4ZwYE&oGYBA&@`7O-Z(&jZ3UCMG6Y2B^G>?3^6ACQ`R>&45)Q-3j7;%77W6~RGsRvkOc)Gvb#b<6}vbnxqDE|!AV1O$POA#QDet#N>MX<8GEIgdS zh)vVmkssno*cbB6=z3~{HwB99u@ol%`dWn`1>(hyqgbX#yu|1dKD0i??!sm;t);?Z7i z^xO%4J!<6jCen=PL%36u$jqDq_YF|?gOHRME3g4o13Cj%#}!HBf>w+ zT)$PM#P@MZ%_seu!LX&}emclgLXp(Mc)Kf&MMmrm{fjo)Eddb>2)1-yjCaq&!&&!^ z=8dxA|LsxVlQ$J&|WlTS6wVrn+{$$+KxssOk3v0)AuX z&R5@aK`A*x2|7?wo5y_Jc9;isHOL$V2kE}u*TfdQa*u0~lpVTkY4+(@`2G8LHGC>e zJ}btU{<2kQCLr*DwR^S8IK&d$&8G%wF94;>76RJU_0%}geo=s7;0i0n;o zbXwGi!Do~mlEYzQVp{(B{yt{l3-_Zb$z#^r4G9VQz{BBC1%&v+2YT#;@X8#^b>rGH z|BLU%-=Bs+x`tot)`mZ*_lR&E@i$<<>oik`uxyLy?p7u1%;}HfHRS}3jXsmQtLwqf zpFxG#rY0s(W;(iz&!3-`xG%Lw0j{1D^$;5$|7dqtQB{>cF`Q-e2>4M%*J|+nUw=KM zq^v5>t){Yz3ku(~1r`r5h(@bpVq;^|)BTSx-ST@2~@S zZLD3zUNF#!FHBqa01|?2eMgX(n(wSDQ9?q(@7K4tpFUCg5}Gx5>5_B;ODGKN>;C?J zU>QV4Uh_^_YX=7hJ-tzrBYKL2Dt31E7=;cQMx*T3im$nTqS5xCO6?n*t~L6e7sKsm# zp2Km}ozUIY_2l%lL-ba_%cP{VTBIngIT()lsG$gkfqO8Y)+Xk>!5R}%>NWc85SWPq z=D=;TY{O@3V!T*bTm)9g!oos!rikxZl?a-Z9r-&kn?;-~z-bfFm+?MB|+H4NF* z4Zi2qEnDr~+cQX}x|=s|GI9WhU0tjG{CAp}hIo1IAHR5W(QM*vx2$6$N_CTin>n49Kgp418l`#9ZLt-P&5>~l z(5j=Ysk{nVnjp+TsQ@WaCt*Rsm8~sPLqmXEY^PB%F$^k3Mn(=ha;8NwyB{BP^!COZ z=#F8-u`=r9SqX&wF22^u+d7L@)aT~P>oxkER&zUU5;1 zz2@QOc9OzplKL8JZM!1~$cd*#o1{8USC0b3Hxqt4s8HVfGB| z!NGx;*FipTD8w|8uwft3Lgg4*9sNyEacUeUGG}`22aQ6dY~htFQU?0^tNWiPCyfGL z);`K7@AY59JMAgB+6u9x#Lrt&mKjQK%P&r&sC z_G7mu2-pJdpwNa?Y7pcyhUuFmk&&4dS_9-)5|$$baH(tJ`tIzh94ad7024wOq;BK2 z;;8ct@Iio~3!^R4e<0zpV`BK!YsHE$?#tR9>!tsw>6Zf&2Xjtw92YCgBf(QX>_XrI zy?Hq{kG!`+qTgBVTtpD)01@*$tg!6R2i=?u=5^WH>((`J=p>FE^nA!;!YK+k8lnKW z!4(GlnbKv3wO^ixCDS=k?@b}px5)0dLp1M+RK=#+t=}Sso9p*qzf~j^K#ZL^rV5Ff z?`7@x;1pNvzenycop?A)8vRb+%uv+LCh;Foco!wlHYmkI3`RN1@MoLLF-*&riXXQoRaN~(I-II9K9JS*2o!s|jS53%Dc^HA1 z%ZjRzQF5{q9PZ}kCP3a~>N!?wD&#nQ;z#=UF|)s4b$$JG)OmJE2`3Lve@~Cvc}ku1 zBqbr?ukG#YGJ1LG7eI)5)hlUK@x{7DTP?@!0P{;#D~`V-@9h!(kzY1jZxILt9UYy$ zgpiO2YUSGCJ{Ei;!G*ei|30%+{^hUpijUFJuNHpZ3H;6@Dq0i|v;aUmrP;1#ATsjv zS;6JwD!;2#e6cv}Y&@C*m)FTX-*?FV)v(rWZpePOt+kbLm+>Lg53FJ!C$tzk>xTn? zcp`X^N^n02NU!K-mFCtTKP>+sj=d*{36@G9K>wY1y4psUkwwKTdP60ESq{8B;(|fk z2}R&N@M{SIe}x^TRkH43(`aa@5LW^+i~lI0EE>bKhe+p6muA(n%NoQ*_~k<&D$;i) zM5SJzVLkSER3a|M!=p(gZ6RNnQ(Z0CQR407;__;6XlV6*(BEjHW!UL{T~G>DC**)c zdXQ{rXh`%-t4LQzNAYa7!GCaYWJH&m9#xTDSZM$1)nHqjf;AL9K5pRomxJDG@*_iBM6+m(QQS8R6ty#6?HZ{=E5h=&@#nbqRF7R`NpbRTTaWlDQ>>n|omY_&uks z`DiSyw<9AXataD6L5dwr?}>GE-mG%r9e|EzVGh7gu?(q zpP88fK*YyzuBNEi-rfBVfDsiD@j6`Bh5HtqO3TO`ri>U&7jMyS{{Fp-Lb-Gg)nm+8 z3>&(j);qoDI(=K+2pVp1NXd;?j5!jjA6pxGpTx22XWta{+|L$)vz|UPHRW{X5Q%cES(D$+j^EB~gK;K1E$!3h!MVAnni_r6Y7P#LVTCe>{{H|i zqALT8*a<~-(|lmtPIwQl;XRM)nJ%dE4)1RbVP$;H_>fU@cqR_`>*eL;TBOO_%g+Ke z7u0lvwC1;fvo)1`4JFom>0tK*%0NoWg1Yu;=d| zKmCB0WbBdR1rnOiTrci|AIdkRK`8Nr%wA~H?`w8fZ>s1>z3dE-2s#{WTU%QoL_)hd zI1~e|lWqVs4o2VJV%vy-v|Epcv?7VMjg1OjKG({~@Nj!)0=Id*XMwV*sj03K%T#Q$ zjvzfjY*klB#~(RJ9W*pFghfEz-PPr?B%?&g<&<#I%!JcqxaWEp#|#RsYGv@0o}|L- zwF2bh`_(kmR8%Ft$-CR=DVr+nXLkI2d={6&%_<38c9L+Qk@4_GurmHd?|?X|3MtPd zu-pXFuk+)bF^O508(;&|S2s18TNuo=cpW2E0*S)+^Lu;M{}m)`_8-{8V0W6^a*B%` zfRZ+@aRgqrHae+%K|^Tvyful`+|F*>#^EN301nF}OKGv-sdZ9;mlH*LuK82IVS(Ln zMq8C=@G_oub9-BlrC`KHkP3xA(+0J7tzYVyQzBVuynHc7#ZYGVJZk0~NV3F&+?k;o z7n?jxQh1Y=W}!0i?5_=PlJuY^u(A96BkH$#ed*XKtAzdGtkqUn3w=Eb1?VjGw^^Od1n$+ ze$b}L|7%x;?bh?iuscvl+P4ocEiC967-W3CKi)44eL^o~0CY~OxL+Kp5ah`-!lQ!& zenG*))9r0Xn??(h<#~2PgQJVi%ra3=-&=1gG3EgcX1b`WtHZ_ou{M-Ph2r+*h?rPg z`=B8a!@=?TG#qw#;id_B8A1C%PEKyB#>vCOBQh*Z$N&CMyPhyy{mY}-npf#w=b=Ya z+V!0`zOe-mdfngJ8qKYFyWGbih13mg18znv4j_Jb5pG}85* zb;R<>!_&W7q?1z9s5#lzBPIIR*FdCk{?aj?Bf}Vk+V~mlsA7y4mnvmfEeKk4IQJuc ziM+<_^UWgv@fFXEj49_IUy!xu`vaeF8nO0e?i-Z~iaLnq)gj48(NbQzlBWKskq!z( z#B@y#o2Tbx)L=~y@ge>t`Djh`{41)aVOs3H{twOfkNgNTCCqL4sLojPJ;vvIlOyoM zm+Ln9?bhEBG;YG0wrte|su05^ey#PEJ67|Bd5jDRa&s#UP2X>vj~cxTRSVcdr)OV-%{!>AROR-zbVdGFg;fTscrPvN0yC1H*E}GuW8g zC)E=*np5_Ue+H8xY<3o?3iGR6u%ifcTkc#H@nWXT`@SkU;gE$`6?!eK*)TY6n^j;> zUh4Gn6XvR4wOj4MR)c;sCO1V%kLk{eW(mB`g?({lnToi;S%}LuVqf7ow)B%|oI9uP zSR1BN2{OIvtt6^RszA~+YoGS46~<%bjy5J4%T%7(NtmC`O!V8;AW*`bf{!I6MrLx- zap9bcy2}Kr#QMRN8Y!x`?hF8$bygtcYBPYTk1T1}&*K!Ano^@?k2&RnHZgNXRVP)s zw+@nG^+p|}%#2o#Rg1K5xRW8Ooh(ioP;~xOhrxnq*0Zg{+vf^-%n$v+;DRo`*&=Us zzMpgaf{cibV=6MQ5GQW>!i+*h_T&}M`F%pQ9&4|kZ?DmKrJ5aoi>>WddGuyj@?;#Y zw5~7(?a3q%@LV@HT!se${s~{zuRYzAA+uJpA|ZlwVHf%H znEH~E5+J$c7i;{V!v8^AE@x6EXjVA6xjll*i~g0h#575#dH*{JE#2KQl4Y}r&$?>v zdt8n`N^qZ(rtd|n2DP}LABsI@m&<2Q0JRk1-8KZhd(Ix4FjYQx4-_+=)o2aW8m7$lD$yZp6PdvLX&>S9;2v>biqCmuI%4Y<$)gE+^3 za-q4jir$`?dQTUb9}ShiK@P(w$;0cHSwYOW!Piq!isXQ zeJm4aMpr_ZIkvd%r*Ekv)YfnWyFe4HKZtNfL8&RE_ot+H= zXm9RC&?6zgU;i(J4Y}_MWdL^hKJ<@oz^z;6=dqF!=caQ?K|*rmi=O$nwgKJB`Xq8^ zTx7NOh`8?aY$`JLn~zv zdLO!CYfL%xOPze2bG|>Fli(}Bp{hrilPCbYA-h$YBJVQ$Rbf$4XhcMOhPg76QUD;^ z?@3U!!->P8D^mIB`BWh9bk#dB0o_BCczSySpfNTXT5xCV6NFDP!YjX~&dTVVI#p#J}2Q;aP4H8)H926bU_VrC$EjBuL2f0vq2 z@kQTEC!obnZFcoH_dj?WjJM%fe`vokss?KCV}>6Q6T<@;>kiIcLg}TMi+?7}F2i~5ITIF%%AqY0?{X=J9L4p#pbZ_Zy)rFm}m$n1b z3^4lw95uZZe?n!eYL6W=I1IRATF-jKv6+B54}63FnvtTbDM}pv$w{f(d@C{NC%x9y9rbw=e#ly66w;jA@COtI@hB7>%dmepxk|>=V!ro%CXwJ^H~!?7lr+ zu)?%n4aEO~teuln&Ap)YJmAxFa&m%w@UO3&c|bzgJYL8h&|@bK`Aj0`ZO9BU9G|1sos zZ`ccAi;Ig#Qj4$8&(jeMZ<7nmMHCJ)V(aQ0A0MlHDnh>4@$<0L%zQ>YGBIKIZ>;*d z%cgx4;BEV_X$U^)(wlpL$Ay>49Xan`5 zEVXgm6W;#ibplV$XPs5W>&CJFb|m_g@Fp-2{`ruo1(9AIiI0Xo+H_`=_xgM0@V!wda4A)i>WQ<9TEgocjzxVTit27~v0%fy@q zzN@@^aCcaEIF+~;KQFKIt5>g-(zJjcP*%PhLIU}cFZFw8=jwDdy(=s|eXcc}VrF_8 zYQ2}P0s4P7J#1~;FBCwLvx4kWT^lDHP(sj-cn0*47JjJ_!u~+Qu%cbV`sGW_{jZIU zX^VFu2Ai}dNOw5s1#xuTPQSXa(e*1athDA$@A&ysR#Q{c(^D9W81e-}DlsaGj_Gr3 zAdwu%j+~m%?$XlDnd0G~Apc2$oMmr7RPG05f`&J6d23*>yA!M~7Z zW5+8VhQw{;s#*Yn_dAj#@ErI04%T#9~2Q}`rm-`5(VMK_cuWxHg}{Uz?1s7 zk2XG#L)?bUMe)AHg%Hz$_ypr0kDoHhNr_~S1ghIqJ<<~eUX|Z4c-@+`1|Ul*{9RH> zj_>^+uGrYpA-u3pg~!xA7xTZ@XnaJiF8ei+QBQtW@bHT3IwK2Qs^-2k@La6Fr%`A+ z9DC-Jzorl}H~o=DLPX&EJq)R@OLR~8%uY}{44OjWq-f+tB3E*g%FNiYe)@#w>we_H z@0MOzm3syhp?mRY%SGt{G3Qm{$v{SOC6O@x!;vC&$6W=HK($zpO&F#UoZq4`J-%fZD|f1W0y>9kB?wEfAxWae07T3vNE+YB#K zzcY3&ZI7Sq9J<-aJ6vmGblx>ADPj#1GQq8W$hl+fAxXsJ7 z8e3js1iZ9+IAxhlcri8}N8&GRRSV57XYg~s$W@K#|HZ&G1rKAkDW!JJ-}c@6iAhm{ z-YN(OazvjmUJVAA_$Bj-8+dVVevEX>xAo%-)mfR}d6RFen@h>zciAQ^MQ>Y)Jd0`n z`)gDMPJBP^?g;ux=4kOZxj%_N_YX_GlC&b{K0AJ|qS3)`xYr6MnGe_y>LcEBDMm-i zorT_-dFT7k}^+C}gOYlJY9xA6Vd0y*&k!F@k7 zDu~a|Y2=*^5l|vUKk(lIw9kT}FyAS2+ zvIv!Qo}ij2&zEQ-{SLNbPv$l5`&~&|%#|ZyF>7PaeUqkbJI&0&41_@{S*O+kTXTJ2 z8XM4?{`%!Md%?+`YBrbaQ%@UCjvr~VD5pJpD{|TsGKXp3P;*x&N!{EXA^I6Q#$-e( za=6o9qDCgf4U@@3tZBllJ3iV=6i6L4ghg&vA9>_S`ageSR=3s=rUCxjGF)@`Y|qyk zdLhRrlWU&qbL8tF_5{tpM-WZVhn|U=6L9?H6%l$jg zKisr-p}h>THuN)nYDUp^b`@6vEMs};9%p1zCMbTnczecWKj-|<>y2k6x~jO^yiFq# zRMVmGv!zY*ZMCFyzt`)Fx^^_hs-7NJ`^Lw)*>k00N1Y98zeCxH^k&li7(j=Cgo8yU6m@U~b~#)q|M5s}<|Zp=0e)?%q`i4}q(svw*y&)DPyi zU-yxh`*XhMW@@>xjj9^hoI0N0_QG-BgT`j-O zKXDXEM^!2A&d8d^_@Q`JI+q`VyCLMNj&XDk=(^YVUXv8ba8fUNl zRnTIckYXjf`~IqGU=mO1%G1NzqPEH%&D&-Vhi&1JQ^x+1+Oa3YR zhlrD_dc!lU#_F3G^F+qx&o^$idaue*79yC`*3SAwMNZmSaM3O7#h-djotMTuM6{@i z5P{1@B_{9UMwQhT4Y;QbEeGBa`_rH@4LQFS{=TS;INf?PB%a5;eDTU=P|`u-sPA6E zRpr5je)3oo$Mj0g1+8&)+(a3PF3mW=d1W#KKWf4Qj}S*#&) zRzy|pZ6i;yZE`-;Qk>3WeW7hNdnGswPfE2l7c426`Q2X9Ft_l2?NZuuwLEDacENS4 zmpBNt51ZmVx{^QbEsLyOR%p|K#_^kN-TYGAp7<2=%?HANL!tgjrT$-^xf+1oJUu-P z!r20-)!@jsw9c;CFwxngn1HMYxbg2FsgdkVgQh zr4#{&!@eefPeF$D-)4dT{@h6a_ppC*V%l3z<*>nuzwy8Dpc(+3JfsjGg;oqch|d3H zkT2H+goHp76nv&mfS!H<5V}9)avtRBWo1rsaxH+aDJv@jTNJRwU$Ro;v$NYfJ7p^% z4h95WxC#Pcx!XKBISC>L0Vke^z+=tL&3#gL&3u}R^vkwE?^q8-xcr@FK)Pbw@TtX4 z9z#wd{&riJK4>m{%=(*XZEkKJpitTY; zCLXY#@NsZZySqGwvjYR=o~RWdG`Zm@FE6iGE-o@MG8E+GRqsg7D+>zPSXmWGg7pFxGTTKrgucse~Z zH!q7f*wnE50{R-{XibC=2IhQS5CFGaZP zC(_5D&jv{6&*-|*k&!LXqp$l2bPP_48hQ?`Kf%Mt#~*DP7zI`pBhda03$&IYRoWW{ zp=dU%@7KLg*TS*(B)rP~GgBhBd81fcnze zs?~y83l|4RY&BW9JWRZ)y_@1lld4gN8!|PQ;WVg-<@5LN0_i3h z85zhUC!>LJ;AvyOXL_Us6!i7cljMXCogmE%js%AmK7Xcln+syV%24uG4qpaA-wq;|Y6rset<7_?K{Mg6{ljR(T4rV@=;j4&EDG}SU}=H4^E|&}Y7p2@Saw^W z5@*PSWFx?mjuRT|qQ@>dv;n~f@Nz)|kQk{acLi8zqdazDIz@v#o+QAUa9KNTc>z{5 zSO#KD&o8OY+$od@1O7!J{ZlFZ*U$ewsr@acBr7XW9n}19bo@d>O=s&x%s8MW)(^d5 z-si$Fii?lmcr*zEO4d~ZU`Fr5pzF7m2fXY_769QmcjL)Io!f?yv9Vgfm)zWy#YjoN z%hOdM_~h`=P!gH;bHR)sEtk6byba&I8)S{1+7Xye_~)^(_+7EqYfgH2FI>|@eFoPy6+L_5@8xQ z^E-nzsHQY+|n5}i&7zJr&aZ+Mp zKHP+))4kNZskDUSMXhyDGHSK2Yn#z=abe-f^}L83(fVwZ)PQPy+Kj{Pv){o2hr>N+ z>cd#15Ps-00E^f2Fo`u^R9Gl+u`6T9I>Ahw-B=u3AKR(npTk2(C$+Ez{G7RJcnUPH zdhg`qHTr!&u=3{`{OvT8%Pm^&PT0xF+`3*?{~Mg-KVQS(`ITfaC+TAY*$i;F7Xnd| MQWm($dY`J*fY4 z&OP@&&;4-k{l+je`?q7QcfIRfYfqqpoFoPs3EH)5*D$1>!IiFEyP5u$G%N>CF&6n)b|b^U`VE(Zkr27;)?q=ia? zy`M9|fsf{5N~>N#(ro_Q6DrJ=qY;LYQKQYWG0pp<6-f8OxRk>Cd?+vCo4=90+nTi1@FJAC>mT{YgThp`lp0H9~IV`ozS9dZAv~`%2fi*Y|%9cr!t$v6K-XKYsMbByZv)B_?*A zYx3LQ-%m-QI_G`zM6+k6Yp`yaM70$yzjo=eQDCoX^&2>jJ0jE2V zB$QTIm?Cg#X$f!cPB0q_3+CfO6-M>vCLeKO0T1Bx1y$5P*`$)=;$YP~E*Pd!pFa8d z`@6fk-oAB<$_`G?fW$+JU%r3;zSOLLK9guYoI+@FVj_@IRbAbf^$P-lSX^2nmX{KJ zvBAy5^A7F!hI_07ZQTLv-z{M zzFycEheSQ2q~uk567)L>*$6+sK-Q$h#K-lO_GOCwAL-RsH1A z+gwTW-SU7!#(=eoeyaBw#QqUZ&|qd=-SoL60|Udr2en*$??WjmsSekT@#6Dyj}5=# zJHBTF0};g#!}sIk<3E128#Oa+kfrf+Nh27nX(U%&yqjujXy}7|*&xj^i+S4Qq}JI{$)RaJ zclJG|!?@OJ&F5crRh*r9FN5_}7M{{9ea}`Vu)6&*l?%^LAAj(RLbj{q&*-&D-+%zJ zM0#Z3_)7^`08=2g_tY_IQha=AqHN0}KK5v3wmXKRaJ)zF9#X>)IIr%YQ$A~CW7Nh_ z6Mc5QKNdYeA*Gg3ZYWfe3@_lpc~TO^lbdUA&|k$Goi|m|ntgNSS4!cs8ryjeubfwZ z@Cnh(CLDyGo}T`UftgvS*j1e78~v-U&8dper#5@P@g#T~&@7W6Sk$+D@)F+Kd`S^q zDEGq+d1@zcRB~aeUdvT&gffLz_>_36Z;DBR_D^-0bG(ukd1~uDt-#;KpFcmXkD$`w z6A-W}rZM?F=HtNK;bLjo)%W5dqVfYIob1V45d&AUVhIzIhfOyg=m&i?vVWvp&&7y* zZD*%&%d)K|z`^WE6s~ddjvmou%WsoS*hjWUm!H zM8y4Me^_965aQJr6&WO$Zw~#+fxGPgQoJSK#_2u?IW7s%4VfCPoV1$iKJPA=qezvy zntcm?aMP2B9b>G)!NE2?E{XFX5`7oRknzb$My-)J$e@L=n$*uG*W8~h67FX7=HvG0&w*F1 z)<3ef6cL9W+I{E@CV+Feo^fcX7Cb)?KU!Qib$wdZCDlh|SV|VTWo+D%~m_}16a5m9(c18;6%0Gs1;YlzwD_;3gUTOF6W z?ZyN@t%;eL@%I7CG9sghc8`r3*S(CG#yK6FAjZHc7@b_gx5Hn*vd)$5-6cZ!vf?+Z zf|-H_k_wj%;N}!B`B`A>`RorLwzLiQ=GK+hjZ(p;>dF~$3N-NO85nSE0+gY*JdE6; zF$Pw#88Q85#H<3eKP;2`yhfxRm$W(4Z(9FxDYJ>=TsdiJ4k<$QST^hcWRS%??a^jS zDA!J8MJe2zji_Zr>jWp@&6{g`!BM+>F24q-Lp;WiAkx>P*UFKw1Y(Q zu{_%Wdn^4k=_x7SYQN&gzsH`JhPGV0{@`cfq)W)RCkNoLL;OkE@gF?(O~`C$nDrR_ zH3jYV>3gU=KX$kTIy!oXDTScRz_ApvsR{ORP&!Lj&%dX?zkhLY+N48!CpZ7s z9rTvAesU6;9Ou=cM<3s`+2aI#SXfwCqMNhhd1FvrQ#18WeY7aL3ep?@9f9xO>jic) zQdDrpU~}@FXb&s5*RMz1+^Y4YHbXd%A3r{34SQ*CpZ<5cI`zn7Edzs;n7Ke|(uMaL zMi|D)=gwgm#iKoD$)0y?O~ElgZ+t`FzvBQ$P*70NleNw_1n}7Dv0`~N1<%aPjI%#_ zbX%lJEYxhz_!=(LLV}^lk<>^hPtH{Ug_a| z@Bn?7uh;*hhld9n|M&0TW0&H z{W$m-J3Bw0yBv=i0sR8u#Vppy-GS1W_=?#P;Nalk;+Amhkx4*OA-W+TiXsrVT*Wv_ z-uLwNy)ZQmtSKUhR;58GAn>Rsx}HCO?&9u#_x9}_DJK{`46sN9^dTUW|1bZU$-%>; zrlJD&b!1|Kfo`VS(JCAV0A`?&kkH)RoUpL)&d!d5g99|P6%OW5Anxsrd`}};TU$#@ zM>j2pFf=p-gom1%nw*^cm64a1mzbE3Jx?VY8(UF*yk+TZ&pWggl@{O zWX()&ZEkJ`Xw}x$J$Ue7Wo1QDLIU>mDKIBbPftMef)qPUbS!ObLc_v%SnmqVZ-}t6 z_Rr2z-p0hj3f>2<(9!WMl9G~Agp<>_L|s)CU+%&_53t0w*Zg~@RJU(3Gd}x0 z*WBjM$x7hbav&1iyOaTiucypm)6nQ`Y;6%V$g{Jv&u5zJe~p&-TpZIO;s^`i#?;k2 z-Dwn`N`6b5-|9(7NXXv8%gby0wX?TZ$YB}%j$21cYN}Yaie9DN-0{xBVKDZBs<`;= zcg!0BX%PXipE-c{K$t~DC0wr_N{9?-3xn$O=ONM2ghNX+O?NJtD2Q7M--g0r0ryk& z&4)vE*_lJFFLRunosGXW4%}Kd_|Y!9y|NPh{{7tg`u(}ikB`|?XfA#>`hw&Wh5-1h zsHj*j4&q1}#_s3y@o|m4_SK-^U_%mHTU&(Nyj-f3rvV-r&Vr$f%dy{m!j?pu;C7C} z@K%&_V-vw4y~gf%Aa#Q@uA-_t7h^s&MxdP&s`B>@tMoC2dR*I)Qg5`wL9bi z3L@VAvSI+V1EvfISiZdtb`3WO&j0g|)Z{u09LDaX7s}W{5!mx#=4bO%P3i9{Ch$;D zUrDpJoQR#8A<*;4Qd?b}F6wJoa4>K`qa#HG08{#$E=n;(_0D;({>a>%A%JC&ZkA~! z#>U3R#B^+Z22*HJ#Je5-6j89tyHDrWxo9cuKP>lzC7oA;t~UBNc0sS&tG zpu&s)K&Y+{L{+uFTCMEv-MhYT-?oq#qh=FP_IVZaBZ0cr#1hOgDK_>VCT98JBp*uO z=wQ^y@mAK-$rou8l_q8Ai5q zcqsS$c|vkBB{A_m42<5M9uTyI-1dJRAM@@04YcR^`S~ClTV66~kBEqY#ZMO%M!<=A&zAWkp?t(kno^u&DdO%*_6ECBeuy zAE1f!u9p@SU9&g<3R6~PWhG$eR_$62o*N9IC@a>^dU+wLrDG#Fy9C$J*T<>}*aQi9 zK>>?~AqaP`x@PL!=g9^@u#7;nfhAI}0-Mw`yBrL3}M)Z}Qo?ZmYSF?O!t$e3f zuav$-Mnz?$r)vs~#;yNnp^CpNzgEV%HG%Q<^nkwhvX}B(f4{Afkx?16*8S)u9Da+Z ztE&r9ed}3Ec*v2Zt!;Q{sCW;+vg>=XW0EQy4@(37mfzVD7 zX#6a?q=6a_I)nMXDw2-HcR z&5-D~FH)ppL=i^((S510$2kw1$fP|uxe{D%+~8y^VN*U?P>d3;U<7&}zgP&qI5 z^74{KG7QMB#3DM10Lcj^{~|dPK=KILp!vxM;7q#Qt%JBf0uuM}<7c(pC9_R`e}8ZY zKy8bN`2*@|X#mvYaeUoiz{!5vk6oY<=vRNWUey&}=}TmK{8+cXLFo|q6&~8{av5pq zlG4)R>lWg0cuK_n{_FAkt#|I+G5RmMdibN=Wukm*Ybz)St$2EP*q3BC8s$CK60a9d zzKu)J zfb@@AYN)6nBHlMxqar0|+G-%!9;=zv{|wpN-*0?#OPAa0(oVFgWyjcruq=jT6;mjK4W!_ZZdnch0` zg0bQqWGR|O)7@F=<;#~ZUSzl+pTb}_F5ORdO|5x(c-q?Ar+(*t(<&wq53oe|`;U!` z?2Lt=wn+_r5+Jm+vP9W^mdCKYKV!(q+~-gSYI99mTGQ~bGU7FwAC!vfnxBE>ctKIo zKpG?;---p6C`YF!-!aR9T8`^r|Ll8UStsQ(-uJCobxV%79zxKec+MLc85x3tf-S)x zqmY}URNn*YfRdD#|HGzO!vUwftZdUYVsupPvI}KHy?#y)wc|Ln&T?{t2l||4Q8o*s z5S?D8mL;x3dBYXI$g(V#*|H53*c{@O}?+>ie zP(0xW=7mw_zVJ*_;lHr``tuWDXxvk%u=}+OFgQR(fT#^p`-Jh~;Rh5HTMG*hK+?v` zTQLcTb87oqT>?~00Q4p82Nqu0K_PuIWv9&(Sb=w$Wt{CY)iA! zG$siAjVh=at5ZRtFoL;n-$d=~ipD7@DUbH{v{j+hCVMG>sezAG-MRW0b<^nRXiN-I zalTqEh)JIZ%{Sytq92rK4&;K>>{4j*m`@NLa55OT?~Z~$k9*drc1cf z6YBU|Hr{W+ZeGM&2bGE8R2QPs3y|Wp9FC92G&CS{aDngNa_nXS6WuE6Rem&W9+pA# z?XbO#y!Y_%5Ct&q?sT-Y0L`1_Q&Up`WPlJL8BRe?NvZE;w{?T-PHb$n!`6?gp&=z} zW@hG%sS5d21%n!p_6i7);!yXk?&y=5_InoSH-y~;1r<;$F^qhAyT^DQ;3n8Km4%Io zQvSrg+MH*NM8T8LKk@%w12}+jdXN^GL%|^&bLV7di>c4?qF+?!TpgL}LEhKO$?;FU z%IWEC&goN8JR*}mCc$SgMK)Onhf)xJ{$#!I8=>4>T=bv1?h~>E2yh<;0R2Gti6+0a zu5MyxCY|&*9wng6nc~F35zj!DF<{QchmC2YGr6(B%yT>_5D>uLLz6J!$b6pn06`RV+yR8jau;#KA?{+-`cb7{0kWBTwBlan ztWs4%PmP-K;fNioAYCKhrK(XLuXEzLoL~o10g!uUZ^+9unip*C5T>{5> zl3;GS#@YVRR!IwxFKcUNqk2$6is`S-NxKBn3dB&+%hEDO2qaN}96P->RtI*-=>4_I za_QS2j?go-elSmw%8Yyb9F73?xBKHZDYOR=uAUx2&r`1Q2e`PuXS<_W3uO_2gbO~& zP+|N>6h{7wD6lXwq0$DE>A~0jAb6E+OmfUZ;&zUXY%MH+YIAaOk|~$?VH>B(n~+rx z?+Eg~cX+t8p~2h9iNF;B!ABFfwJo%_puc;j_6GrP)QF1HLg5kqY?!pAHq3Z_9k2ed z1<6{ki)vR@91iFfJgKu&>bo)|2JBvhqe=Fh%_@hn{tvZNrRnMEB0rEUe!o#6_7Dh5 z#Kgp4s@O>}+-TUZsQ+gkh*cX$w1t3B%UH|H%DTF^1PFMvLGyBR_YM#1?HAgD+8D1` zgXjx{TN>HN98iQ)0l-ctD0mKvUGt5cD7YNu^kQ_^0<7`ykl-@xP-F(kM2`w2uAq~G z=qpIv0EI8pc4eu!fC_qFo`3sx&6^dK*wv&YH{86>XzTGe(F5FopGlTD-}BQyuJCgD zF@Pa zT(&*_jpuU%z;xYzlAEsGZIF!6$^1_g=HmPh3d7>LOuzmj{regR013dIUZ;Pko;d>?llR@g=5>U5!Kbz zz2_sRXJ?@BG{3aO{p1M)U0NT2BOt-q+uJKED_r^bK^>1mpdI|VARaK!xP9*dZonwp$km{vXy5jxb=#bq4Gdi9bc?E=`>%CN3YunrY6%iU58ZL89ww!F+XfHBh%GG!m#!AhiKiY2%wG%1TN__k+aA0bVP7*I#gyX+Lz_e2x{Z=Ldsd?0C}9-%x8ohU&GPR*49>1!o_|vc?5t zV?LPrJpuA0MOw6wz`*x1_aP7>a`MFFE;_P%<(ybTqVAsok~UsXz&NIl}-PfV^B)Ttive7qT&wE;C<3eLnc(2YTXjd3ga` z91_y*tNiJP*4xcRc-H>bbj>*ApAOgk^XE^X3)@eX)YkscsH6YGqht|4{^gUFl=RO< zz5vh|fdNxJp#kA>=Ui6%spJ;UjexsGZ$w%)xIuH3yxU7V6~%_pTYc5Gqt0ck#8fwHZL_R_yn3BM895 zNU1g26lga{&&VJN9!M*`(p}iOrVYVHJgynXD;E=GD0r^~N*B~7GQ}#$4DcQA8~=!* z=4m7#0UI`dA;n*;CgunNugp^~X1^Lp;QwCmiO6S$g-QT1w8jIe85)_6Q3Ic`x<#Pj%UIY{QF=`e|u|-g^}@e@uPgyc0?6e#H!jSK?jj7Mnp?Jb=Tt zUfwn}4bDm@3% z-$SLbPyUHf8|Y+-6B+e4!z?#< z`yLUEuZ!?L9ZtS+<3@pAU7YPLu$~f^14WCjtQRlJ1PG=Fp=8Y7 zdliZOm%~3sbgq{5queH5Dci&s7Zx)3SKiWTuBXh*nc1nndEcqwG*7Hb6mE?4g&BJK`^nSlHxU-Lk@m8SF#!uMVSh7Y9OxQPeJYXeF8QyGH03xsl|!6%92K6qx$)OC2VtP2 zxA3ItEi}gx1`9&ZK|L#~)X~|wJDdpEeabtq6z3V&yR!P4nwp?F<2vjInd)XPU;~mR z>NN|j;f%G;N&dI*UY^gisq!97FNt`Y`pG+%l-04{F*5O|NJ+cA^g6y12ugN@EIM{@ zLr`Yj1Bam7qm~z=0(CCMpyefSw>wh8pe4E-dtKjo#M( z3Z~M+GAR%G-Om z4<4ivIN)@}v;;3h5$-KcR#v%?Oa%i_M$pj`75xVK7+Mnux_$Z&x2A)uHJQ!4yeLSs zyg`L~zWb^NY~0!2zBoHusxLf-6^Kcd8pOe^fNlT$X`rAm&|Iyv!G)8?8XO#aVkQrH z&;D0+EpTf87A%8h78ha+%>(Jo;%iXbx39ul@N#hAQf`qoW0+_vH5X;mOPG~dYlbQ{k6GS_84Ygh`IYl91RT(`&*zpqFkJQ3&$ur zn2=Cl6zt7)$n0!g|^YT&vJq8pnCn+xOZQlse$w%u434)|}4D`gIDnO$;2mm;= zvV)V8(g1?XF>|ygc*6AN1*$m$CoiVxb`pm({0cr=z&hP3nWXr%_Je#^YWKx@0sL2UUicSzR6%x0d`>;eNnT~t#E>5xJf{E zzW1WKP7llT*cb$k;~lx;0rxMGv^ zYkwvUUp;yP_DdDa3sk3?`-Emro-<^0K2PB5I-~&sy`@!%Z;T3%*we_w6H<3aR*D#F zoMDqw_lY-Gg(@D9&;#04P42WT^4c}%*)fWPC6aL0cm;pil}E0P1`?JpSSIsoj-W9 zSGb5_I0bQ_gEf%QwDW&;8$=vP@FuZ7E^u1EduS5Iq(phFOp@7&9c1SSCP69eG2Yz` zN||{zHKYunaB@}Dw=8o5Ew&_~ZAQ?V1n~V~a#}!gYaL8Y*P@-_3@>YUY5$A>T0fl~ z9UCn{TkBm+G7$4+Wo766y00Pq_9CUzTj3wm-ue4m-J_N>%{tD+HdLRO!FR~evGMlS z$6y}c>5^Z|$w_$i`drb$gic!@(RS5FW+ zhRwS;a7di!pkPoMe*{#UK~{Nke6kR7NX^&qEJr~H%=Vo7LimY~H|i_158 zdX7a6>!oi7b-jXjKvDvP57ZUb*4BXj)8blWY-eXDD;o-S8R!b`hj|Jau0THXlviEZ z2Gr=2hkI*6tSSk+?2C>)R%>1`Mi{79040r;wGcFW%*?Q^XdB-|;)kQNFf)%odx*%R zoP2L>(Ph4w%e3slD|fQQ6f1;-QmETBr(q03-PN31KN4Al&eOC|ld zWS4uBU;TX?hOAixAlU0Kt%zS)a3iuHxtbj?ZQyh^L4c4~!ydHm$g-vuoVy!MmElEd?Co_v6K^jPZ;Xqi~0 zCD+x(90hFNHa-6+apmi%O%|(3VLmC}%%p8w!2pNbK=^l*GL6%D&(%5^Q+d{P)zRHA zGvDq0ZSm%_au6KKD~dU!I5Lk9A$>G)MQ#gy)FQPMvppG!=^4hx|ADwJjIi;v>BR>OCG0EAa?w&rX^CS?%xS{aW~{GEY)8 z_I?OzcLVo_|Gswv)ZX-kSq+!)AGAGO3i&E>^t^qq|5pa>#Oo^<6?~VD(^^LNW#y4u znwnaM*OTS`z@Y_SJqm07E04H=N@ezNN2WG9JxX=k|Jo!u_HF5Y86Eeg6p|u3Z#@uJ#;atS3O)`w4%YqgntQ zPg?&hrjjFW#uz;LcKFmuEh44TIt!fGq~3ArOQ&}>&rn)hZkoQi@}kHc97{lW%ZyLA zG!VWR?zM@;F_w|XMaowePLDpxnjC#Rf1$d`%eDhBC3aHF45+(#bt5b5yfl{rHGZgb z#DGm#Ygf~2oqxILmx1&p*IvIr+HhH2YB@AyXSg_wbUo?UC!W3-|NdRb2?J#o>tQd` zes!-rdZBjoyZN|wZdV()R{1ffTy|%Ex3G-&38_XA!_WBPg~SaK)IN@q`@WD++{XY5 z?n@bw$~C0D&_928pWiqxM6E`}Z9Y?hN_a#Oc@t?mb_EqUsCp$hu!Eat$I+P}=Nx_mT-O>te5=#?hWv;O|1Fl7BDl;dZ$@Z4C{xdo5MC zr;1u+{Gt>dK*!+I@Gl7a;C^2}F2^VA(b zpf7Tq3g=hmBrkkA9r~+P7fC^|v46hyOZ4WS2Zc7JsDefaH@&n=S@)8D%#axuq%jU; z=DfN6{?pH4>ZyWDW0N9C0ga;N0cZUC4>lg4wt)xJF}@vmlet>b%21p~!ENfFrAUNb z5iP^m&gZA@`2G!aFVSHB8<+~ozy?vbg3byEh$+XHt$y8Gh zne!P27d}XSUV)x^R4aP2iFUqPxG6prqs_b{wP4+BiJ}}G0Lil6GC;-cD#eX!P2_{=FWOj zp269X;`WnVs(K>!{mmpOcF)v~Kd&eMp=nkI0Y+a=TkqE+aG({w$@9?3t0la-_E`4@ zD1`ta=@2?%YYg=`HQfV!~c3y^t4_rLy;F4#}L+7D-p%n1tf*CZ3?w>-rPd ze1--fIc>CP^tL>j>W_kiW9E|W3;keGwiWdnuQ~Va#QsQU|AW4qo@0$O-EE{SbDFNY zOF-QfcdizHvvyB*n|+6(k(JeF<~S^a>H9HmmRP28*-R`yxybRVzb>mF&AGpPcXz+r zwjdJ7PC%>bb9=6kt9v@H_ErP#ZhkcDIQ!Uy@``Ct-1xrMRsf*qc>U49+xF-!kCf!> z$kMDccQc35cNa%a$bHF@bN*!SjJ5pPq)FEeUm<+fTXGE^aK(YeU&e2XbVu-KFy|IN^jFZKMw;JZnUO7Gb7NCY&xRoInwEvr^JnXpUkRj^W7-p^pBCRys9{VH4ksQO#7VA2xO!8>GJD;wjlnVtL||7_3K}9 zHuf68Q5?mP=7S!N039H&Fg^^gzoQmqF5XH~oXk~aQPpEe1w3m_I+-2H$vldbLa zn7cMV=lp%i;uLlg*?DKQsH$wcM1SJTljDzv(U)#hg}|B7^v-Nrw!^)PGUwF?_v$rs zdKhJ=g{I$H;_vDI?!&gj@!u1XG& zi$Pn+A@9UM#LFBVfe-nc!Nad_F>}}F=1deYBa_mOJMRjb?R-26h|J$_I5I;P%&dIO zNUk&YSSmd?6Og-<2?%fJZuq_mnDD1Weu3F&0bT(?E(QN^lln+~n)0UzGb3TRkNyU- zoQ*Q=>ET1r^d1k2^gw$l77`ZTT3G?hw+fu$;w*|g{#du^Nth_KGUODbOcejPNOW3`b<2-rC0_~Vqm^=yT3Or!W*ht{P9hTHEde%VT6Jdq6rpXYOPpI}kv!9aG~QZ;?d9HHU)F|v zf__|V9aQ~Z&A(4c>izP?HlrM}`(RPTx+nADYCfryZNZvy+`feCQ1BzVWY)k7DXER{ zuK7#-9R`t|X>B?rPI7!PQeAwbB1g_k*TZtYK`Z{aa#J3tP?3~p@&qEL$8fKM)6_zJ zk>3YKNL#$KbIZG0a<=|y(Ve6eJEDtxK!kt2K(j-M>7c-*PZ>FjwkIVKT_eUNE;8@Yr0 zaWv)Kvf^wjc)(hStkAFD^=>$BLi+YQPv5&rR`F6t0jdZkBgDbPzyP^=;3g=n zfU?UWFY2l+ZmM1(xcXF(ovoOy0(2^~{v=ax)W>OX-hzr@KmcTiY_6&4##uZzEWkYw zRB6H03DhN9RKfN?fAZb+ZTu!-aD)THsU`#^M$jk@x{*MuLe(Sr=g;x*@KQ|f68dl& zf|li0P`DECeB=XKvOq>UF`=2s%OMR)SWc;1+L`snM>ajl6x4Wgj6@U^GY#I}p!8^B z^2K4yj)aH^0)ce&^$9y|jL$Xuv!6dn9|z5KDnzXkl9KK1?E)3tp!Nw0n1m!GYx|A+ z9q(zt<=cWG&-L%l8}*?l!!$szAO|{Jz3blc?uDrg7KI=m@PfP#s5@PyUd5phDmFGd zYwMmJqlNi-(E2Fu2O5#siujAU;WUfO%lyu4e$Utz4WFhfeVS?4w0*CW`9G$-_Douu zk%2+8-rM03=guo+>|2{AFCou3wYbb7E3diTPxQ`px=m zBEThB;{ifn1_{v+H@CW(Mh>$VFR%?PxJYghfGeIr+C8=fGIqvhz2lmI(`?)MD%k1* z%yS{ooVL8UNCg3Vc|0G?$`)YoRVjA_GwR3a=<@cqt%oXxr2RSW0JyoC3wrN((xN~& zXYpoU4KvW9IWc_BPknw4Qbj4-eCRrU?4sR(c)+8EoXcgLicrsRIhx2SrI|MStnueTn@x+$w=OZKrlu0zA+E2B_?{Ok z7lS^L`23~_~}$sZgu6{YY<-%AKUMcd7#mv9;?2pX)PBOfqr;kBOCSl$iPX0jw(VowiVG z?Z(@F$x;)Z+W^~^O)n_WoT@%d;8`zXu*3q1bGXvoM!r!^YS?wx4V#xZC#_`lUH z0AC<=^_8IoVG%EH_Gu%FI!FB5(AvgFAhCN%FuI*;@ZX>Gzi|1#LG-^rQIrsd;#t&S zw*L?R>=-q*vx^871z!t@=iuZlHEzcYl>n{sC<1tOA&%!duyWgmi$fXKy2XWs&kFPr zphaWvbiEKX6aVIII0@Y;Z|wnXTi}*aT9)_4nD=8wMh=$q-KFkvP!Zk){VAZ8%gO1m zUF4E|2(&$WTf@b~uH%C42yf)sEDa}Src(AeH)_gx0swDV8_r)TD-VMrPdlZz2k3OB zhIQP-j89JHohgs0`toHl-20q{2lQ`JJxf<&AOx4_Fmn%%j+(tGJ@*H}MRyaNF3?nd zxu)Kr^kb!LOz)`$C_#ac1lp^Nwo$(sQL>davIfFSF}Th_jEDCSejKcS4tq#qtKC!k z_3N!EPz%+mdi98k>RU?*YyI*3!TQ+J?>ejEXTsdC8U7k3;T9k=U_D2xOUA#Hv%r1Rp| z;EwF}Oq~D?O|u5V)zx*o{;Z%@MoOv^G?YxBE;u5r=>x{SF4pzKL5~6k2G|wJ`N<4! z9IYO7^gB0i1;+G$`5hD#R59yCdIBm%IOpIFVK)>$&s13Z>S#;%38%@)UNW+_vXaP$ z1}bBJ^TXBb>c1I+JC>*)xp+oMPG0A*B1a1S8l(AhZcb5C(?ZV+lm+(>ce`0>Jit$0 zlt4gm3}r=t1{%;$!+i{n-hQgwhSKYpdS31j6P2(Pqz;vrm+N2dMr-9q<%(63h^FCEDP|%C)cFJ<=_P{gSY<&w@}=Z literal 0 HcmV?d00001 diff --git a/img/scan.png b/img/scan.png new file mode 100644 index 0000000000000000000000000000000000000000..f8ab885abfc5626a6937445df9f80181788b580d GIT binary patch literal 19180 zcmch9Wn7ePx3A5J)BuuFLx-SzO zeV%up_niG@|F)lIWahrFxniwwHvC}u3fi<+yMV$ z=|44l?V8jlNfDvvP9y8}%g;1bk{8D>o9c4#Nq#~}aPDu*$O!a2V>yXmPL8)0kv*Jc zR;jLwna`|KsTHfpTXz^`juV1GgVAoZG@xSf=tI9T-@4%wtsjM@&&VVw;T7dBD@bJ! zDM|C4gKJ}z&H0HWd4$ByX&`aXm+6T~KIcNtl+)V75*p>GdTha4#kE; zC~zq#{+C{o%Dxq&yu{fO#DdeoF`L`kXnA=xqpc?Dmgnbx{rvf|zInXXF+Dx~KB{k7 zLcP=sg8}9(3PhDZgoB;^)mRk;)%5BrF%C|;Qm(=>9<0uO6)$0YaFF)=VAQIyHQ)gO zfvS>{(z9o?sXLN3EB*9)2aiv6OkBw!6fJxrRSgZK#KiX`MMOmCVmfAK^xyU#J!W8# z%Kk~>4jqY1NeMg1`?aFQNDqfY(jb*9D=WGIrKP1^U9yTk-Z==|aYY0@J-x{pMsTOH z5&}74uQIc;nz5cedo~bbotBc4QdCrAAfDcP)WtFRUci$S8x<9m;o(CqEv=XV9aU92 z1_o3NjIUND-M{iw_Fd}tMmytJUgE1}fBc9#3bSx?JNx2~nPicsocEB4Y47k*g${RX zcxfr89AY(A^=7KRzMgKf!F_bKZLHQ&zC!8bT~}9EL`1}la7=WxswZ8qMcM95Yq08d zbMD5NdQ8`YHMP~)me$tc)BbL|<@^I{(b&cP;{@$vD2T8z7Qx!h06)+E2Ity#Q! zRd}RXQc_}RX=(In_&|^V?G6i9am@bZW0Y|T2t}AYVkSY)+^_zVpI&!l9D1GzuOb;znN`>cF7uL@Uj7n%deP%YD=t zA_Qjm`1p>Nq<+sW0=l9iBMUXFBO@ZNleoCLb_QX6*-PSaY&m)O@FCqx{L&bG@~`hg zJNEYW5)z1dE%GZG8eWaBef})SNkG<^p8n;yV9kA~u#lPU+PZ+_#sk#INXIFIxDAs; zi_BStFJHbKJ*K5?a~aI}_WgUb<$W1H6NvBSoe$EPop%OMZd~d$96yFpwh2WuEmaZU zrI_0ew6?H7`@(i-wWFgWKAv#czJOc)5xD zIq3d+Zn#fFLPBzKi;5=8oe6|XXx_qpjn%uVpe0ny7G0~Ts7E!^xQm&Vm76R1QHby< z#AMupuE|B&D@udwqPSL@$i7tCu}rq2&xQ_iSL$o^Z8^>sD+=i|yygna7zr{mUMF%b zDlBF&d$D|2*kPo6s1S(P+owb!QGb>N_7ICQp?x>pY5kGDh{!Fa>!$g**=LU*KPGTC zyjf9RPJw$-PMS(fLz8hoD?3}!2kJank*2p*rkqz^SmtPICl7lY7AC$!HLnUgMSKXa z9Wp<-s3<)zKCKe3GLIX{8+(e|ZIz$R2ey%#Yj1+L_l^AHamPuu#ennk^A9TPfmrVL z+cTnPQU*D44wMuKUS()dubHmyHA?#KvByTvuCB9tI^R>hcDK@fYSpe8pJZlcZc%?# zSoRCItA2s@W9v=X#mo%M?6s{eNu|`<$U*fxpWq)B>@SC6>ci|W&YWGlCMK0D$OZ_Z z8|BjLFE)szoYdJcbBVdn`ZA?UWnGT8DAYF=7yWOMTxvM9_a#GG&zjH9&iL$+gLws2 z$5i5M!+$~tmQ>+WO$|fg{_v=QogFUa+VCgiBCXni@p1Rd3-`HbW++9S?=NP}s;Sq| zAk6#N*nPS_K1Q2XdF}z$WLjoLuLl(bWhlVR&oq6oYCIp%{?cz|Z)8OC+AiNOT|re9 z;>e9PJvT?CEEc9FuskcnfemkeiA~{`?RvN&>6o04j==GJYuElQ9G}~xL6O(d%>C($ zvFop2h!GGHE`wXP5eW+mldO$b44;mgk-^5qe^0(b?(+|rHZ?x84Xd@dsWOoJwO@LL ztO7J5D^PCiO9*f46>KVS4~HPIcHW5u{B&MLk?ZQL9j2Anh3|M`1$ot+8EP;*JS z2PWKg-batsIoq}gv+0NV`1q8fd?+i!eQIxSU(xHZE(V7<%fN2U2>JQ>-C3o@w?++X z;-;b!c|I^WxJl4kQ(LP@Iq`zVNEz(4U>y+K-t#zOJniGH&>Hu8L!j}EU4HGPBZiR& zJHaPdK}1ekI?AZIY0uZ!7o1^?y3()$3(=*`%}q|Sk>9ftEhuAk?dHv!2kO+xEal02 zk#TXfeEgqg1^4#%r^VN_R&stXKR*;kFET)_U84N_`Ey0ZH`8~GSF`0)V!w)@DcRWA zsGrT)o>eZWXE(yRgzLb>w7DD!^tZeRTW;0AO{uY9` zq@?u>;&2XKtTH*tf?l)o0P+c@n<*5RkjVuv7%tg^+(-N@{J38a^YfYLZAVHku@-#u^a3IO074Ld9diyR0|ev|C1g_M=WoHMQfo9o@EV1`7Zn$$9IuRv#(_gb zodWFa?K88prCF#b!=mQr=Rpd#!B;*gH$#3ca`yY#SzS?)m9_OF<5C;|4*)Phz^K6A z{)b;Oa`W(fX>N9NcYmg}d7^MH^LoEmQ11sQeqk#zB}19eFVg@NbWkPuM-6>x)m zeQ&qjd@8y`l$DirjMZE6xSP?UH_9vBl7v%GBj*2vl~B}E6UlUO#%>4Ua z1_lb9m6ercW>WX#-oO8po&7`VFEjZh?(&KXzf!4826}p->wQ6?p<-AWD8=~AFk?Y3 zUfwYm#|piY&_VWrt|*cm`exy@w6yESu$e{)INUp@g`$PxF;Pgj%8q;jR;fz2N`Zj; zt6$qngS8a3>nke^(YBitd~e}kaB<99-cID%8yl~ZL8XJB7}F<#Krx! zi;4;gTK(F8)zs9Ke)?2d!J>Y0bfk+Dt@!ij&rb~v9sr}@7JGmL$1bFb76#x4gnC_? zJ>OdI*|VR8$SHTBp`r2i@rjFzQ`}fzzwb|GVPP?KH0{NuQ%%Vld4d5CK=BedpoPIX z>8Ytjqbb18pa?)=h%~FKucsiU6*Z})n2|OlT$wzLEs zIH|qyDutEd4C~!BTj9nx>js?sEoPTU$jrdHJ@m7cXAW z9q#RAWM*=0<@H+(7lRvuwS~v&vBm>T{DL9Io*LtBH_yoEXzxDfci6mR1GpV2?TnFz zj!sIGZcaf#7l2By_Vr8{SXkm-?H?Zl?n4Mj~@@(Un_C}u%F4AIcgoX!c} z(^*>T42mT)8=;t$tCy7S#(><#Nk~fiRWL(=PM6aOGP8Yn$O7>1?yk>PhQ9y>L4OL5 zV_`;yNPxA82^~B8fI6a-I5aG5R=h=CPA)d%*! z^Kx$K5j#5`6q=?$mr-lo5V$!H0j;nSxfKO#)P%Vl=oIm4DK@W)Q>I=9_#_!gi04 z1he!WO0{^vcgYS&1a_b+BJz@c6M}5E&Mh*V_(fs1$m5h|g<7 z!%UEty1GQU`1nW!lwZ6^8~gb)DmHfEb5UPg96BQ-oQ;x!;p3-IyN}6qkYmSy+C=3S z74_&o>K_;|GBT2?Tza$tf~|DrlAD{`*RNk=?VoU(9r0Wi0DhR1d|H!M%Sx7IdM^gj@&b6ghzdgPDuEu zluLT~_yclvtih4@3&KUcjfUbWNN{m+0VV~6Qmo{?A2cmqC$@{-fYd^u80hG&U%ztW zO%ELSu&a`7M@2=omepyA(v4ml<%CK5Kw3$7Ld0xO_Dl)1&ZNeO5JV-v=H@c4rt|*v z=g(e)VM0bS;P|&cZb^Rh=lHvNqiaDl26@IY z4MCn+Eoz(?8nXU*5cUU8pe~O(3W0d9s`?JV9Vw|GvGg=BCx{0cLBYX{3=C*@zAP+c z<-{i>0J=v&LL!P>)b{pN@R^o}kGq}|klo_e{vK9iSecZZ3}&dHu+R@X{81D7hYPTG zU@rjK0cZq9xwo?erw|Zm1mrW{VYx3eHWm*zHEL8 z^nmMxPW=gfXk5s)g)0m}I(0L>S5;a0PRZT7O)zU8xol}kNx+;d@${^$SQr!^C3b;* z0K)U{GrqpO{D#%`^=rVrEe~(x**iK?k{?7yM_Wk(b6&-M*N?+gzQudc3s z_l{t=zCGxNwyG;kQ(e7To0Qkd+8~StnX788Kf59-U=K)0&{0tn-oIyymkIYTfanQ1 z-TtbssfjC)1@M4L+P-!wM+6l1MOsECTB*{@>zaWP9RO?$F;Z1{^Lu-CXR0q=P}kxk zkLy3e1RdFAI&<^z@GxL@21|fvL^D{8*BH8wcvF^?yVK&0y92%FLA!VuY2Z^s#=8`4 z*F+;Oo~L=(HMzBn(laa`n`KGxg}S5k_!jx2|u-?@2$j-#-`{PCBmtrv$K^CZqtK_4L8)%A}M2dwXnaxuQ$)*xRE+GGO{u|OR2;? z`g=Mwpzz(?++2p3l>oBzbY{5SV+=79lN@_>Ee(y}iArsB=9c9bDk}IGGB!4kL;dTe z;O#aNz27x@l6Blq#y%=HE{DPpTvAd})}f(K_Df4yZ92Qu5#+SAi(f;C`@&|jNHT8o zz+xT{5tTMP3#dOjJ;levJ2^hqyS<#-yS;65mZLxy85NaT8%cmo0k>(BOYeQp@2>8G zX5Xf5@EiT0+__WdA>-*OAX+8ArPHR45ET`j(foRnGqooby)AIotZ?042?KD!v{y>uGigyq^F zrNpP`@IHR5mw49XRvN=Qhj?&vB0$Fn5_S>A zih%LrLKC`-wi#}jL*?Y;K9HLk85`fcaYG@zfC^E?^RqLtcuHE)d<}J`notkOtl)?A zn$+1Y$|)#d?eGpgrlQKIuh%LXWQ>v8F1$4?1rMlAUm5j++!u9mxV^0jm6MgFrlO+5 zMeUbxq^i;<+a z8~@`46J0R}GK_C?hC^T^XLMR^h2-Egw&2j$8OBSI=@{#(V1c|h6L@rgl>+{PZ1ENj z@@+&NF8H%R{dZyn5*Rv>rr(Qov~0xd%-y?sPi0>po4e=H`d&&N{0}idIY8OUdGO%Vq6)g3 z{+{YHO-)TjMZHG^>>Q1rRA&DH5<0=wzks`>7zw<|-TRg9C|oP`fTf4oXT&6C}Xi7I~e(^lbSVQS#^KCt-<^ zm%?8YL*u+rLJmX%^hjo_dQF%Z8GVKjypK*^$_y^s0jLbl{QzujeO+{Z;6P;woIVy7 z7ND`G@VQoGXQOm{|Mo4npupaU1%feZZ)c~=XmO*QybWM2Zj5}O?1flvYiy4%l-m3o z`wXyZ6GkU89aX;qkE>SV02gV}C7}%cD@i!PbndHW^;M$rzj3rAQ$Q4Z(;Xfj28$TK zd?{RW-LDRnKD8W@MmMMQ1F3`@IoRYVo+Q{9C;|b!d+l~%WQF%LpiUVZ85s;lhulm7 zB4|2csg1I-vVvF552v}f8P|>Chy1>l<^^A*+W!S7b>9;r35fjMF{JakOVAh0!7tR*SzRLZ|E{b=INn zYB76lE*&J4Qk(ynN(MUqkxKF#Bg2q;9t+bLtj?sI8|$-WOsB9P%Dlja0cse4Tr^R4 zeHxG{<>*?L+TSTo71%L4fG!}FiQO&s7YhAWp^W5W z9YI)K0YO14Yiqo3XQjcZ9rOCg-1+<@yu-fIygJ|M4hl-dnBePs_wE7Q-b_6JVl8c<8BJq&D^0|Wp*v2m*(qf6L9I*lA#!IUb50L$xJ%%}R zHdj}_bU1E|300Su|7vR!&#A?C;+YwcohzvCdnz{SR9;GX>>c>++bk_LF1 zgq$3nBemoGw8lnHJw0Fqk+DZ76P1Thma+vXjFtQ}VMEG}q{idD>c8P|^P4x&5C>YY*{Br{pFv}=bSKYjws{WF!22zky5=OgS&Qcgq9F8X43B1WX)@h z@|nr~9^vi}LiUuqTCLtUTz#%acr%c^uPLmLD<=8sZN^qC__|w?YomEBklo74o=1y| zi@UqIv9mX%-~jf(&&f$wV9K88=yq;lY`my=Ve<0jODGhI>I*-S)AC~X_w_wKImyb% zFkvG?Qk{b^of0sBkNk(pQj?RDWH>mK9q8!j*>Xup6QP4U3B#~nyHOZ1b8m9oduC=P zCvpNdLji&v%ug!2;~cpK_+8wF7_j-m@&fdSNc7_rX?PFu_t1O4d#w-f7Z$%@c z#(_vxSs9|9*J_yFy`p{&R*H5WwlOv$JR0K7M3~c6D*->+5@Pw0|mZMVp@w z>3y4%lS6q?#|RE8*NVwL03!mf9>1(L8{FB2xbEBxInsh?^ZiioMb$ zl@T;i^iQ9@$W*Ex>_r}6@z4W35PRf5h4nl?0iJ;=J0Rh+Pa@%nkTUgxb>8ODN-OC` z5+b6YkdSh`e}z6BpiJ-e)8z+}-A5ID{koK`9`l9|#8>|*5jq=Fvi|tVp@xqFWd$HA zSJMiJ3*Z6U(38_sB#0V*!X*pf4L$lD2Va05HMF&p%OQp#q9B5_+f3!%00W8@YA`=V z30DX*&$!aHy;wYb8c!C%Vc_TzBB^Y-$MD2o=yz~`lLSs5wHWcQDfs}G6R^- z_`gFrc^YIru(+{N$K^b)Nq1{!hxJNp565{OZAw@slX>}3um|AQG%e`C!9h^m-8Ug2 z=u5KMKaryMPJjP%6pFcaYj#gG7_Ul>gR{X|iERsQ8#^Up?tRf51??Q&u3908P3!(w zaP<@hVi6h!2HrfNab8|tS7TLE z;ZLu8Q4{t=peX{u7?dk0<>KnFF1vvlYVYF!Gw!+bAyi9fObMaTZJgrncO$a6sNqB^ zO+jM#ow;aL`FV|9f50#>C>)Anz+sUztp(9H;ujdx*fRu6-BvR&TEa2-565itMpOOH zr6D%SpOwJ+CXh~a(r(Z<59)Zj1J5KaZMu`CacsncXtyO0XH5qluv-U3V&`z%)z42=N zA-(SS%$l6G3vZiIFD!B3@8@P^I%_Lj`^9Y-+zl?j=x8_2I~^ODbx$n(z!Y;>y3RmE zOzdN&08WpNva$mfRNru@*@O+&ohcnB5*w&xXJ<#=n_p1is~-OHTcIflit_5l#@C9Y z0w>})i}#?aL)GoSReo0woxl0rXSOgiD?OBq5Ti7-!-RV5e4_>OBj?ZDa&OO31eWn% zIHP$RhvKV^fgFJ5Ahi&jL@+enA|xbK2rht7JIn=Fz!0t?)!LAkm_nNb}6uqto8L_ zoSa9+EcmrDh#vma{Pjd$AlniCrlaQpp%f^%dakZKtl5kzg#n%J1c#gBwJ7LJ%*>=z zl+iM_pcp|y8~dhebO6F8C!3kxmlJqrs)=Dw8w zRF7L-u#Xk-DOl%g9XEG2IpX+QP+H9?o%W{ z@892dN=~$ohuC=UcIsU0_+auM;WHjB%0g`JH0pyb^1{8%fq^gI8dP&A`mF$y=Aj*o zfQ-zM9KJnzWh$+!B4)Hn=B@|#|N2Go2Imk$609&r*+PSKh)L5d9`qmnCI7r#!#fU> zSpLCIRrUYUmwS3`;vti+oAm(!0U$$FYBUL0!2_QBA^2_g=%*lTL zIidC52NL5e@@9M;w}?ZKOiy*~nsy4;2lWSs?qir-nNYUQs;*ym=rn2I?+rB@J(4cf z?ozmB76RBlQS@$7Y7;^gJwrq<>SQN+w2)|Z@l)XZXP$uHrQ=X_G0gS-L7SM`z{vRs<(0d%5eSD+SAcDl;D1H1E@luzj>}2JXS1jS1 zDaefSrjLsMP=E#okjbnF+5Vyg$zkCAVE(Jp-+$-)E3^J3i@eYRA(kR$^^XVxz2Z?* z&q3k@v?lTot(pGwKI~5dbn*%RjqnYB0E&`#qrp7!>OECv(>18k#N1zB?|BUXt(ipth~zx|ECqa!zwthX_K)FVLOi^M)z!qv$Pc|H$Y1NP&uYcz z>Tw^|zrtt8P#+c+Hg^*fss|#QsHks~?$FrT1)aRCEZFjILqk6REyKvb067rMOsi%v zusP8~r@?~^sh@GI3r=b^jWx7TU~MX^sC+?fk55HWB*evi4rDEmO8EInv9V#wev7(1 zim%Mgzi4gBg#acCjQU81H0~rhTE=Ez zT!C^k4J|EMDXALw(+Z!rSiQRCz{3~$!d%SHpy=lUQsthDJK)_FqDUSPlay3n?%sU48f7L6;(-zo~-n@q4f?CUqSsS*AZ@a4XoIj&VH zH7M-mGdf(X`>aps^KU|>dZ(TMTTT_VTUtT_E5_&(q1y6aDm(D$`6HjtB>u(`VQ&YMCgaVsIPmf}6Ap`OTz%gLYl-(MY zj2U7%2!#U#*nb}n?~a;&qXakbKNKbb=s$mldE=iN2~uts2}R27&SNqx8US75&a(rV zWESax0nw5|Z-fF^Tn23F+Tz(imw%h>Q82`Wh*$s`3-IZr#zEc4f8*GH)>c)6ejo&! z0bm*NMWh25-WKw)XQ$9oG{RUmb218&#+@&~&R$Vh2lx@a$k5=Re}<)vjUu3oO1$Jg zx1i3%dw{%UDd!c0P;WrtyuTx8)gsE>w9>*tjKI{Yc>i|Q2jbCsw4>i17qa z?1lOHKpx*8m>zY`y9+;lC;(_(4k>Z|BZR~y6OoXSwS}Sj-nbJpa`TN4VkR_6EFD>U zYDcUzMkF$O<=E8+l9lRNigE&-=cV$z;=5BJ5JE^f$c)t|%E!|eD*I|9fz>|M2-YudotgekH~KWc_WHlj7;N(nWo&pd5B0KXRu>Rv(LYcR4TRZmmkDv4Osoi>%2;iAtb<&=%m+fUy;; zgR&by3hGNrdPheY8yv8w#5v`&e|}TmfINNr^p_BX%dIeHqshB5L~W$4LTiwil>f4k zygAuhQ!z?2?+BC7>LP9}Zgg|PffaHw#5EH&O9?jzwMC#h)*?#aoqR>4|NC%S_=`B8 zJ4wug@yMs}+`KoVvpsu8xi9}o#}^W%;Wj>lIuKg)p*3f|c~b(JiBKmkU331b6hmER z^sKCCjNoKq`aU?w0V3jGg8ZMVtLe4Ae+lR3<_`M9A#4aKTidCrj^Q-Lq6OxmYLB*!Q=RaDR{<)rK`^9`SDn9$ur^Bm5 zz)#J1OG*nOk83Zh2>@j;nkKUPh>f%Sl25rra#qdopEfp}p)wu4S|$#wKL2O_oUtIi zd+pZCVwd{&H=!dNHwq(@-vjTwUFx5X!xyE$9EY9V-3XnaVCx~g5l}$~N?AOxiHR9q zZB+kutY>kN)d~+xm59)=^cK46WiwMyw<(Tt41b~JodIPdq$#f94=|Kmq^b#^xG!-rjObZ-BfJ8`s~34*?V@VC1C zwc{573VDtW4Ch%*sTbpUSkp@J^E)GcrsvDvIK@-nb}Gjz$quOCIT))zL;$67U=dKE z2mAXJ!@zv+H!IJ<#cU@R5^1_^ZJu_R(ayEbo9fs?;#zxe& z{4+aqG=64)ocp?lteo6Pq9BZTKp6q3HII=osR0I3hvU-xV<=If#>IS;CB?KcYdj{?wF16MQ%doKd7<0VbHX1(N@)(hD{d;0ZY%E zS-yDdC|@S?Rd?PM2g#oPce3aD3;E^d?g7{{(9?cv&=d9E1=Auzx$<)|6_V$iztgS& z3LSk1_)M&^!Gtp*OtY};d)NkH-l^mQF_zx!xf?QP-KBUH_y=!2d$tUcbn(;>?WR?N zjLU4tKTlooYTB49c6ryh z@%Wwl@0$Tt--D-jr*6r(x^|45ZPjkFe*eQfVLXHM#i8N;re0I_j`qZ(f=|TAiMNq< z>CalEY^WsPfwwWTWY9Kw&t5vb{%KItn{w~o1=VgyatXHs<<-xX)Y|;|BrP^B_)cu$ zri)_#qb$;sU(ZwrP)_5NvoDNJKbV01^w($QqUL?gRF#O{pY&x)ZpuesbhjppAEyU7 z>NaKHliT`0cLkOFzt!DWGIdTQ>TCMQoUNThlN>?lysu)kE(KPc=_oF~2vg?irzck< zeNuMPC9t5@V7GM~P(F&8bTT2-ZZ%igBZLSuV^}21^-MTt5P3H?)7gEsj$QAxC)@;D zS#`F`hsyn*7|~WHG-6REM+hB{D|P1ed$=Fpcc8iYU1F06!V8>9PR`1iHI+@G2I9!% z2>R7fpWBxVtM>bD_@_}wrsw#0*I~#SsmY>|L>jT1vP8%>5+4@jzi+VVT=9^M#@3)q1iEx^COwc`Vx;k8h?y z$wzsU#AuTL-aQzlh<3X+HF78N0X6ErXe^3I3SZ;LmC>Qq1>3t3U`sX}N$^eXtEk1= z73cGW)=MF`m}{!qIIa$z4a{hL^zlz^4jr;J{GN7#2I7Wd`fVk zYm*F1wpMS*F}!+cw|FY!Fgt>y0df11X0fGj1@|--%x^el$i8ctxeNONiTLDYfs$GVz2!(NAD-5nQ`M@P{fffyM^i0~gov&eXc^f+OSARr^bEcRA+6x4=_gC1lk` zRkBr^%naWMS?21&p;r#V)YuUK%ud{(m7ZjN z_d*j4xlXNY7#3@rd?OVW2{NCibu{{#o}^r870!Ri=_$uNq<m&qU`|mAFg3}6P;@lhSk=HpkR^8fsEx#1P zc&aeP@z8Kn=OI3-NLqeobFtpk?|a7xx`t0)Or$^58(1d&^*m^>Ex>T^u(6~y6p1Pyb8$%9M@arGq9)amgJW&6+K!b^cTq$NF z{oeP3`na}KEv{vXvvU<&@0DK&Ebms&HEvG0C`QHTzb~m{60_3GJ+YnnR?2bta=^TJ zO2?>y|Ix2H%-a(yIn#Q-SAnU>GPZ)aRmQpGeSn(uM~>8q_^v^e(?$KBU_$7O(7R<@ zV}P3lj*8A6Za?(Vch1^sDb(S}`NbUFZ&Hv!W4eJoL#*-Pt$p@!S-^Jqz0Z*nn=_Afo@%K;6k;i!RLg zpH6Izhv#1=ceql^%s*uttT&${K|TA(u-Q`|!mE9?Q_(V9B7UvX>JMZ{r+R+>Ik;!R z>3kOiW4|pikJGWkI6=?-{{DD@f(mvbbaRfAPyD$~tJ!Krs>hr2#j7cN?@euKc4nGI z9As(kBNIk@+Lra#_s^}cCMj8B!}{zy*k{Lro?SdKSBcH*@L+>(`BcUQ?G7%|SnVx( zC>?JAax&a*bDu(%vYK>{ zUYV4a+kjV27dQKMt>#k#@6KILNe!RR~z3f2T56*0jr$c*yhR$ zEpvu%g`FY!9cL>YnTH=4eb0)|YSOs2%t*lmHPe1>Z zx2D9@bYvHHTIep3mLwE(r_j;me&N1ZS}p!a%ZmY%y62Kj6uisc!#kEH~C1dJFv0)q{(gI!kmmr_`o4MJ!?m7y49RZ z=ml4bHoeQadcn)+A@82BwL7VE*sig0oUv`8RXmt!cNu7?(Cxe2?ROmAU(JMNOh!ZA@3OViP zmA?PV_dcmxtz>9eH;<2hv6=ie1WnR(Ov}((Bh68I!au~hhz;*z$7p$SJpY)bA~H(r zmG-%3F@OKiyGsp`cT3e>msSXvod{^yg2S5+L5N3|fnZy{V z>Pj$WUQCI&e)R*5q*#R{QkAIthPRAM>fy<6nXJcecbZcU>T#*hbIJ7% zsaZBosRd4KOEiN{dLxPVogFfMVPBYT9LtvqdzKj>E-4`q8((^@b*g=@nC~GyzI>+6nG%y)uFf3wrMcwJx0PV4fvkG9l<-e$_w%*Tpbc>w z7pcR5TC?CCPznGl*TOi!!w)dA(?BQR-Mgk%R)eeiQkgrTx|??b^aJrqAOs7PLFrMj zGeW%}wk;709$*{@>WRQ30zl0^vYPbomsOx8U)JG97$pU6ns{oj2@w$y5N2d(qB?zg zdU^(dy>)bb@}_ZTbMp-t2L}h)*V9{|9tTJspydLSa%X>aaC}37bi#(WU-;WBh~%X_ zpiUR`yF6rO25o*e$F&Kd;ta4XX;7_$g9B`U#}m~IJD}}2A*3Dn8o;Dx-i!rRprAk# z^ix#3p@jqlEG;kF6bU;!wQUz;QilQ9|7~M#VF5G&Sy(n!?2T1};!ay;wB-2sz@Ki4 zU+Fy+6n?C&;C=H zpx!k1`Nx8F4SMio0x$-ZLd{+iHs-s@5fQzh_VnBdxFJB3k9#W)6e?igd~n-e`88Up zby@xt6#Wrwv=Rip!SqQQw&2(O4!U77ltGiwu?VQ{2R(uwpvIkV#(!rkboM1@G;b(K zowZ(#+Mu8c)M~xo2WxbB+~fH*{7uO>-A502kASi*DCo5x!?xB=t>Ox*X~9~E5Pp-g zv@D>j-Ae8#Hy>C&D z1Aw6CKubpCM3>(#vf1q{Q<9RT+ZIgE0lqVs{I*B_?Ml=a??8EtL;k0*_JOb)VA!C4 zCcU_(Cg{yepnXks-2u(h$H&LS#Kd5&KtcCuLq!GUj^m`eBY0SaX%z}syHoQ|me%;B z9VD2=Ra=}Jkc|;x)6=n#ywKUusY1R`o`m%g0ZEXUovlj)-=G<{5zw?qARg<3-bPFIMDKFxyXlVv|2!>D3iinbqQE(Kc82o^N4)qYG)hS}P-fZoFs;3{hVvUt01 z!YP^%`=X$(nd#{_uiMbn zzO)mw_J!LZpBd1UzSi8XFxJaB;07aLB8wdX+B{B6)j$kTMqadqHf z+~L7>%*=WeRj>srAMlgdeik-m*=Y((ub*P(3dL|ao$`zXo5A66ybLG>wzWz)Y5Y5Z~&-zw+Uy# zw81*_^6zW9$%Mii8tOx~r7#lgwl%v%egy93gTZx_0%U@3VM8`#H#}d1wJaR>WqZ&b z+U<|;?Bp~=lIU*WL$-3034C)-fxpR!Yq36u5Z}!cvzT*u?aMun{si|N63`4;#i7Y>IqW-`Z0poC zY{7n7WmpGrTr?HO|vFKVy;K8L~ru zXLdnh@sd>Anmw_LzYuxpBPy=`<3@k{_gzMx!pxDQ;eUl;f3f+`ON1o~CtbS!^VT?1 z|K&3O>f8jt_kXluJQN?47OvewcAxV`gXdNZYP)usumLOlufQKpDJU-9R=e47I%eCG z#ARz_X^H+6RJ#{7UYr2A&Ev3kQZ3F9?r~B+zbo%;H3V8mtV<3dzqlBg5?t>D`)1~wG7Ni2h+K>qW+XDF~}Ej8 #include "testing_helpers.hpp" -const int SIZE = 1 << 8; // feel free to change the size of array +const int SIZE = 1 << 28; // feel free to change the size of array const int NPOT = SIZE - 3; // Non-Power-Of-Two int *a = new int[SIZE]; int *b = new int[SIZE]; From af89b15969a5d3f24f3a3e4fb6faa4730db27ce3 Mon Sep 17 00:00:00 2001 From: Zichuan Date: Tue, 18 Sep 2018 23:52:36 -0400 Subject: [PATCH 5/5] readme 1.1 --- README.md | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/README.md b/README.md index 9b61c8b..b93c779 100644 --- a/README.md +++ b/README.md @@ -20,19 +20,32 @@ CUDA Stream Compaction ### Block size analysis +We fix array size as 2^21 and change the block size. + ![block_size](img/block_size.png) +As we can see, as long as the block size is not 32, it makes little differences when we increase the block size. + ### Array Size Analysis on Scan +We fix block size as 1024 and change the array size. + ![scan](img/scan.png) +As we can see, CPU is of course the slowest. We can also see that my own implementation is still much slower than +Thrust implementation. I think this is because our own code is still not efficient and hardware-exploiting enough. + ### Array Size Analysis on Compaction +We fix block size as 1024 and change the array size. + ![compaction](img/compaction.png) +As we can see, CPU with scan is the slowest. I think that scan brings overhead to CPU, thus, if we are using CPU, we'd rather not use scan at all. + ## Output -Array size 2^26, block size 1024 +Array size 2^28, block size 1024 ```shell