From b4b1afe6b7e23ac22b41885af295ca18e6a34c6a Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Fri, 11 Sep 2015 17:41:29 -0400 Subject: [PATCH 1/8] part 1 to 4 --- src/main.cpp | 11 ++++---- stream_compaction/cpu.cu | 49 +++++++++++++++++++++++++++++----- stream_compaction/efficient.cu | 49 +++++++++++++++++++++++++++++++--- stream_compaction/naive.cu | 36 ++++++++++++++++++++++--- stream_compaction/thrust.cu | 15 ++++++++++- 5 files changed, 142 insertions(+), 18 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 675da35..611fd14 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -39,7 +39,7 @@ int main(int argc, char* argv[]) { StreamCompaction::CPU::scan(NPOT, c, a); printArray(NPOT, b, true); printCmpResult(NPOT, b, c); - + zeroArray(SIZE, c); printDesc("naive scan, power-of-two"); StreamCompaction::Naive::scan(SIZE, c, a); @@ -63,7 +63,7 @@ int main(int argc, char* argv[]) { StreamCompaction::Efficient::scan(NPOT, c, a); //printArray(NPOT, c, true); printCmpResult(NPOT, b, c); - + zeroArray(SIZE, c); printDesc("thrust scan, power-of-two"); StreamCompaction::Thrust::scan(SIZE, c, a); @@ -82,13 +82,12 @@ int main(int argc, char* argv[]) { printf("*****************************\n"); // Compaction tests - + /* genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; printArray(SIZE, a, true); int count, expectedCount, expectedNPOT; - zeroArray(SIZE, b); printDesc("cpu compact without scan, power-of-two"); count = StreamCompaction::CPU::compactWithoutScan(SIZE, b, a); @@ -108,7 +107,8 @@ int main(int argc, char* argv[]) { count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); - + */ + /* zeroArray(SIZE, c); printDesc("work-efficient compact, power-of-two"); count = StreamCompaction::Efficient::compact(SIZE, c, a); @@ -120,4 +120,5 @@ int main(int argc, char* argv[]) { count = StreamCompaction::Efficient::compact(NPOT, c, a); //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); + */ } diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index e600c29..3f4fb58 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -1,15 +1,19 @@ #include #include "cpu.h" +#include namespace StreamCompaction { namespace CPU { /** - * CPU scan (prefix sum). + * CPU scan (exclusive prefix sum). */ void scan(int n, int *odata, const int *idata) { // TODO - printf("TODO\n"); + odata[0] = 0; + for (int i = 1; i < n; i++){ + odata[i] = odata[i - 1] + idata[i - 1]; + } } /** @@ -18,8 +22,16 @@ void scan(int n, int *odata, const int *idata) { * @returns the number of elements remaining after compaction. */ int compactWithoutScan(int n, int *odata, const int *idata) { - // TODO - return -1; + int count = 0; + int j = 0; + for (int i = 0; i < n; i++){ + if (idata[i] != 0){ + count++; + odata[j] = idata[i]; + j++; + } + } + return count; } /** @@ -28,8 +40,33 @@ int compactWithoutScan(int n, int *odata, const int *idata) { * @returns the number of elements remaining after compaction. */ int compactWithScan(int n, int *odata, const int *idata) { - // TODO - return -1; + int *c; + c = (int *)malloc(n * sizeof(int)); + for (int i = 0; i < n; i++){ + if (idata[i] != 0){ + c[i] = 1; + } + else { + c[i] = 0; + } + } + int *d, *e; + d = (int *)malloc(n * sizeof(int)); + scan(n, d, c); + for (int i = 0; i < n; i++){ + if (c[i] == 1){ + odata[d[i]] = idata[i]; + } + } + free(c); + free(d); + int count = 0; + for (int i = 0; i < n; i++){ + if (odata[i] != 0){ + count++; + } + } + return count; } } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index b2f739b..b90becd 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -6,14 +6,57 @@ namespace StreamCompaction { namespace Efficient { -// TODO: __global__ + __global__ void scanUp(int d, int *idata){ + int k = threadIdx.x; + if (k % (int)pow((double)2, (double)(d + 1)) == 0){ + idata[k - 1 + (int)pow((double)2, (double)(d + 1))] += idata[k - 1 + (int)pow((double)2, (double)d)]; + } + } + + __global__ void scanDown(int d, int *idata){ + int k = threadIdx.x; + if (k % (int)pow((double)2, (double)(d + 1)) == 0){ + int t = idata[k - 1 + (int)pow((double)2, (double)d)]; + idata[k - 1 + (int)pow((double)2, (double)d)] = idata[k - 1 + (int)pow((double)2, (double)(d + 1))]; + idata[k - 1 + (int)pow((double)2, (double)(d + 1))] += t; + } + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - // TODO - printf("TODO\n"); + // Padding + int m = (int)pow((double)2, (double)ilog2ceil(n)); + int *pidata; + pidata = (int*)malloc(m*sizeof(int)); + for (int i = 0; i < n; i++){ + pidata[i] = idata[i]; + } + if (m > n){ + for (int i = n; i < m; i++){ + pidata[i] = 0; + } + } + int *dev_pidata; + cudaMalloc((void **)&dev_pidata, m*sizeof(int)); + cudaMemcpy(dev_pidata, pidata, m*sizeof(int), cudaMemcpyHostToDevice); + // Scan + for (int d = 0; d < ilog2ceil(m); d++){ + scanUp << <1, m>> >(d, dev_pidata); + } + cudaMemcpy(pidata, dev_pidata, m*sizeof(int), cudaMemcpyDeviceToHost); + pidata[m - 1] = 0; + cudaMemcpy(dev_pidata, pidata, m*sizeof(int), cudaMemcpyHostToDevice); + for (int d = ilog2ceil(m)-1; d >=0; d--){ + scanDown<<<1, m>>>(d, dev_pidata); + } + cudaMemcpy(pidata, dev_pidata, m*sizeof(int), cudaMemcpyDeviceToHost); + for (int i = 0; i < n; i++){ + odata[i] = pidata[i]; + } + cudaFree(dev_pidata); + free(pidata); } /** diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 3d86b60..dae90af 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -2,18 +2,48 @@ #include #include "common.h" #include "naive.h" +#include namespace StreamCompaction { namespace Naive { -// TODO: __global__ + __global__ void scanCol(int d, int *idata){ + int k = threadIdx.x; + if (k >= (int)pow((double)2, (double)(d-1))){ + idata[k] = idata[k - (int)pow((double)2, (double)(d - 1))] + idata[k]; + } + } /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - // TODO - printf("TODO\n"); + // Padding + int m = (int)pow((double)2, (double)ilog2ceil(n)); + int *pidata; + pidata = (int*)malloc(m*sizeof(int)); + for (int i = 0; i < n; i++){ + pidata[i] = idata[i]; + } + if (m > n){ + for (int i = n; i < m; i++){ + pidata[i] = 0; + } + } + int *dev_pidata; + cudaMalloc((void **)&dev_pidata, m*sizeof(int)); + cudaMemcpy(dev_pidata, pidata, m*sizeof(int), cudaMemcpyHostToDevice); + // Scan + for (int d = 1; d <= ilog2ceil(m); d++){ + scanCol<<<1, m>>>(d, dev_pidata); + } + cudaMemcpy(pidata, dev_pidata, m*sizeof(int), cudaMemcpyDeviceToHost); + odata[0] = 0; + for (int i = 1; i < n; i++){ + odata[i] = pidata[i-1]; + } + cudaFree(dev_pidata); + free(pidata); } } diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index d8dbb32..1775000 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -13,9 +13,22 @@ namespace Thrust { * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ void scan(int n, int *odata, const int *idata) { - // TODO use `thrust::exclusive_scan` + // `thrust::exclusive_scan` // example: for device_vectors dv_in and dv_out: // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + thrust::host_vector hs_in(n); + thrust::host_vector hs_out(n); + thrust::device_vector dv_in(n); + thrust::device_vector dv_out(n); + for (int i = 0; i < n; i++){ + hs_in[i] = idata[i]; + } + dv_in = hs_in; + thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + hs_out = dv_out; + for (int i = 0; i < n; i++){ + odata[i] = hs_out[i]; + } } } From 7ac4fface07499b2defbe61f1057d16f49f37c85 Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Fri, 11 Sep 2015 18:09:23 -0400 Subject: [PATCH 2/8] compaction --- src/main.cpp | 5 +-- stream_compaction/efficient.cu | 56 ++++++++++++++++++++++++++++++++-- 2 files changed, 55 insertions(+), 6 deletions(-) diff --git a/src/main.cpp b/src/main.cpp index 611fd14..1703f8e 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -82,7 +82,6 @@ int main(int argc, char* argv[]) { printf("*****************************\n"); // Compaction tests - /* genArray(SIZE - 1, a, 4); // Leave a 0 at the end to test that edge case a[SIZE - 1] = 0; printArray(SIZE, a, true); @@ -107,8 +106,7 @@ int main(int argc, char* argv[]) { count = StreamCompaction::CPU::compactWithScan(SIZE, c, a); printArray(count, c, true); printCmpLenResult(count, expectedCount, b, c); - */ - /* + zeroArray(SIZE, c); printDesc("work-efficient compact, power-of-two"); count = StreamCompaction::Efficient::compact(SIZE, c, a); @@ -120,5 +118,4 @@ int main(int argc, char* argv[]) { count = StreamCompaction::Efficient::compact(NPOT, c, a); //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); - */ } diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index b90becd..c8f0757 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -22,6 +22,23 @@ namespace Efficient { } } + __global__ void filter(int *odata, int *idata){ + int k = threadIdx.x; + if (idata[k] == 0){ + odata[k] = 0; + } + else { + odata[k] = 1; + } + } + + __global__ void scatter(int *odata, int *idata, int *filter, int *idx){ + int k = threadIdx.x; + if (filter[k] == 1){ + odata[idx[k]] = idata[k]; + } + } + /** * Performs prefix-sum (aka scan) on idata, storing the result into odata. */ @@ -69,8 +86,43 @@ void scan(int n, int *odata, const int *idata) { * @returns The number of elements remaining after compaction. */ int compact(int n, int *odata, const int *idata) { - // TODO - return -1; + int *f; + int *dev_idata; + cudaMalloc((void**)&f, n * sizeof(int)); + cudaMalloc((void**)&dev_idata, n * sizeof(int)); + cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice); + filter << <1, n >> >(f, dev_idata); + + int *hs_idx; + int *hs_f; + hs_idx = (int *)malloc(n * sizeof(int)); + hs_f = (int *)malloc(n*sizeof(int)); + cudaMemcpy(hs_f, f, n * sizeof(int), cudaMemcpyDeviceToHost); + scan(n, hs_idx, hs_f); + + int *idx; + int *dv_out; + cudaMalloc((void**)&idx, n * sizeof(int)); + cudaMalloc((void**)&dv_out, n * sizeof(int)); + cudaMemcpy(idx, hs_idx, n * sizeof(int), cudaMemcpyHostToDevice); + scatter << <1, n >> >(dv_out, dev_idata, f, idx); + + cudaMemcpy(odata, dv_out, n * sizeof(int), cudaMemcpyDeviceToHost); + + cudaFree(f); + cudaFree(dev_idata); + cudaFree(idx); + cudaFree(dv_out); + free(hs_idx); + free(hs_f); + + int count = 0; + for (int i = 0; i < n; i++){ + if (odata[i] != 0){ + count++; + } + } + return count; } } From ffdb00b1fa2c546d5b7748d8280e9aa4636a4df5 Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Fri, 11 Sep 2015 23:20:48 -0400 Subject: [PATCH 3/8] radix sort --- src/main.cpp | 25 ++++++++++ stream_compaction/CMakeLists.txt | 2 + stream_compaction/efficient.cu | 8 +-- stream_compaction/naive.cu | 2 +- stream_compaction/radix.cu | 86 ++++++++++++++++++++++++++++++++ stream_compaction/radix.h | 5 ++ 6 files changed, 123 insertions(+), 5 deletions(-) create mode 100644 stream_compaction/radix.cu create mode 100644 stream_compaction/radix.h diff --git a/src/main.cpp b/src/main.cpp index 1703f8e..e6b33d4 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -11,6 +11,9 @@ #include #include #include +#include +#include // std::sort +#include // std::vector #include "testing_helpers.hpp" int main(int argc, char* argv[]) { @@ -118,4 +121,26 @@ int main(int argc, char* argv[]) { count = StreamCompaction::Efficient::compact(NPOT, c, a); //printArray(count, c, true); printCmpLenResult(count, expectedNPOT, b, c); + + printf("\n"); + printf("*****************************\n"); + printf("** RADIX SORT TEST **\n"); + printf("*****************************\n"); + + zeroArray(SIZE, c); + count = RadixSort::sort(SIZE, c, a, 8); + + std::vector s(a, a + SIZE); + std::sort(s.begin(), s.end()); + + int d[SIZE]; + for (int i = 0; i < SIZE; i++){ + d[i] = s[i]; + } + + printf("Radix sort:\n"); + printArray(count, c, true); + printf("Std sort:\n"); + printArray(count, d, true); + printCmpLenResult(count, count, d, c); } diff --git a/stream_compaction/CMakeLists.txt b/stream_compaction/CMakeLists.txt index cdbef77..bcc484e 100644 --- a/stream_compaction/CMakeLists.txt +++ b/stream_compaction/CMakeLists.txt @@ -9,6 +9,8 @@ set(SOURCE_FILES "efficient.cu" "thrust.h" "thrust.cu" + "radix.h" + "radix.cu" ) cuda_add_library(stream_compaction diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index c8f0757..539d98f 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -7,14 +7,14 @@ namespace StreamCompaction { namespace Efficient { __global__ void scanUp(int d, int *idata){ - int k = threadIdx.x; + int k = blockIdx.x*blockDim.x + threadIdx.x; if (k % (int)pow((double)2, (double)(d + 1)) == 0){ idata[k - 1 + (int)pow((double)2, (double)(d + 1))] += idata[k - 1 + (int)pow((double)2, (double)d)]; } } __global__ void scanDown(int d, int *idata){ - int k = threadIdx.x; + int k = blockIdx.x*blockDim.x + threadIdx.x; if (k % (int)pow((double)2, (double)(d + 1)) == 0){ int t = idata[k - 1 + (int)pow((double)2, (double)d)]; idata[k - 1 + (int)pow((double)2, (double)d)] = idata[k - 1 + (int)pow((double)2, (double)(d + 1))]; @@ -23,7 +23,7 @@ namespace Efficient { } __global__ void filter(int *odata, int *idata){ - int k = threadIdx.x; + int k = blockIdx.x*blockDim.x + threadIdx.x; if (idata[k] == 0){ odata[k] = 0; } @@ -33,7 +33,7 @@ namespace Efficient { } __global__ void scatter(int *odata, int *idata, int *filter, int *idx){ - int k = threadIdx.x; + int k = blockIdx.x*blockDim.x + threadIdx.x; if (filter[k] == 1){ odata[idx[k]] = idata[k]; } diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index dae90af..853b719 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -8,7 +8,7 @@ namespace StreamCompaction { namespace Naive { __global__ void scanCol(int d, int *idata){ - int k = threadIdx.x; + int k = blockIdx.x*blockDim.x + threadIdx.x; if (k >= (int)pow((double)2, (double)(d-1))){ idata[k] = idata[k - (int)pow((double)2, (double)(d - 1))] + idata[k]; } diff --git a/stream_compaction/radix.cu b/stream_compaction/radix.cu new file mode 100644 index 0000000..26386c6 --- /dev/null +++ b/stream_compaction/radix.cu @@ -0,0 +1,86 @@ +#include +#include +#include "common.h" +#include "radix.h" +#include + +namespace RadixSort { + + __global__ void getTotalFalse(int *oB, int *oE, const int *idata, const int currentPass, const int n){ + int k = (blockIdx.x*blockDim.x + threadIdx.x) % n; + + int digit = 0; + int dec = idata[k]; + for (int i = 0; i <= currentPass; i++){ + digit = dec % 2; + dec = dec / 2; + } + + if (digit == 0){ + oE[k] = 1; + oB[k] = 0; + } else { + oB[k] = 1; + oE[k] = 0; + } + } + + __global__ void getT(int *oT, const int *iF, const int totalFalses, const int n){ + int k = (blockIdx.x*blockDim.x + threadIdx.x) % n; + oT[k] = k - iF[k] + totalFalses; + } + + __global__ void rearrange(int *odata, int *idata, int *oB, int *oT, int *oF, const int n){ + int k = (blockIdx.x*blockDim.x + threadIdx.x) % n; + int d = oB[k] == 1 ? oT[k] : oF[k]; + odata[d] = idata[k]; + } + +int sort(int n, int *odata, const int *idata, const int passes) { + int blockSize = 64; + int gridSize = ceil(n / blockSize); + + int *dv_in, *dv_out; + cudaMalloc((void**)&dv_in, n*sizeof(int)); + cudaMalloc((void**)&dv_out, n*sizeof(int)); + cudaMemcpy(dv_in, idata, n*sizeof(int), cudaMemcpyHostToDevice); + + int *dv_b, *dv_e, *dv_t, *dv_f; + cudaMalloc((void**)&dv_b, n*sizeof(int)); + cudaMalloc((void**)&dv_e, n*sizeof(int)); + cudaMalloc((void**)&dv_t, n*sizeof(int)); + cudaMalloc((void**)&dv_f, n*sizeof(int)); + + int *hs_e, *hs_f; + hs_e = (int *)malloc(n*sizeof(int)); + hs_f = (int *)malloc(n*sizeof(int)); + for (int p = 0; p >>(dv_b, dv_e, dv_in, p, n); + + cudaMemcpy(hs_e, dv_e, n*sizeof(int), cudaMemcpyDeviceToHost); + StreamCompaction::Thrust::scan(n, hs_f, hs_e); + cudaMemcpy(dv_f, hs_f, n*sizeof(int), cudaMemcpyHostToDevice); + + int totalFalse = hs_e[n - 1] + hs_f[n - 1]; + + getT<<<1, n >>>(dv_t, dv_f, totalFalse, n); + + rearrange << <1, n >> >(dv_out, dv_in, dv_b, dv_t, dv_f, n); + cudaMemcpy(dv_in, dv_out, n*sizeof(int), cudaMemcpyDeviceToDevice); + } + + cudaMemcpy(odata, dv_out, n*sizeof(int), cudaMemcpyDeviceToHost); + + cudaFree(dv_in); + cudaFree(dv_out); + cudaFree(dv_b); + cudaFree(dv_e); + cudaFree(dv_t); + cudaFree(dv_f); + free(hs_e); + free(hs_f); + + return n; +} + +} diff --git a/stream_compaction/radix.h b/stream_compaction/radix.h new file mode 100644 index 0000000..cb828a7 --- /dev/null +++ b/stream_compaction/radix.h @@ -0,0 +1,5 @@ +#pragma once + +namespace RadixSort { + int sort(int n, int *odata, const int *idata, const int passes); +} From cd4649ab3b11b75a571ea91d479397cfcd2392e1 Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Fri, 11 Sep 2015 23:22:15 -0400 Subject: [PATCH 4/8] Update README.md --- README.md | 156 +----------------------------------------------------- 1 file changed, 2 insertions(+), 154 deletions(-) diff --git a/README.md b/README.md index a82ea0f..3db0b8d 100644 --- a/README.md +++ b/README.md @@ -3,160 +3,8 @@ CUDA Stream Compaction **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2** -* (TODO) YOUR NAME HERE -* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab) - -### (TODO: Your README) - -Include analysis, etc. (Remember, this is public, so don't put -anything here that you don't want to share with the world.) - -Instructions (delete me) -======================== - -This is due Sunday, September 13 at midnight. - -**Summary:** In this project, you'll implement GPU stream compaction in CUDA, -from scratch. This algorithm is widely used, and will be important for -accelerating your path tracer project. - -Your stream compaction implementations in this project will simply remove `0`s -from an array of `int`s. In the path tracer, you will remove terminated paths -from an array of rays. - -In addition to being useful for your path tracer, this project is meant to -reorient your algorithmic thinking to the way of the GPU. On GPUs, many -algorithms can benefit from massive parallelism and, in particular, data -parallelism: executing the same code many times simultaneously with different -data. - -You'll implement a few different versions of the *Scan* (*Prefix Sum*) -algorithm. First, you'll implement a CPU version of the algorithm to reinforce -your understanding. Then, you'll write a few GPU implementations: "naive" and -"work-efficient." Finally, you'll use some of these to implement GPU stream -compaction. - -**Algorithm overview & details:** There are two primary references for details -on the implementation of scan and stream compaction. - -* The [slides on Parallel Algorithms](https://github.com/CIS565-Fall-2015/cis565-fall-2015.github.io/raw/master/lectures/2-Parallel-Algorithms.pptx) - for Scan, Stream Compaction, and Work-Efficient Parallel Scan. -* GPU Gems 3, Chapter 39 - [Parallel Prefix Sum (Scan) with CUDA](http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html). - -Your GPU stream compaction implementation will live inside of the -`stream_compaction` subproject. This way, you will be able to easily copy it -over for use in your GPU path tracer. - - -## Part 0: The Usual - -This project (and all other CUDA projects in this course) requires an NVIDIA -graphics card with CUDA capability. Any card with Compute Capability 2.0 -(`sm_20`) or greater will work. Check your GPU on this -[compatibility table](https://developer.nvidia.com/cuda-gpus). -If you do not have a personal machine with these specs, you may use those -computers in the Moore 100B/C which have supported GPUs. - -**HOWEVER**: If you need to use the lab computer for your development, you will -not presently be able to do GPU performance profiling. This will be very -important for debugging performance bottlenecks in your program. - -### Useful existing code - -* `stream_compaction/common.h` - * `checkCUDAError` macro: checks for CUDA errors and exits if there were any. - * `ilog2ceil(x)`: computes the ceiling of log2(x), as an integer. -* `main.cpp` - * Some testing code for your implementations. - - -## Part 1: CPU Scan & Stream Compaction - -This stream compaction method will remove `0`s from an array of `int`s. - -In `stream_compaction/cpu.cu`, implement: - -* `StreamCompaction::CPU::scan`: compute an exclusive prefix sum. -* `StreamCompaction::CPU::compactWithoutScan`: stream compaction without using - the `scan` function. -* `StreamCompaction::CPU::compactWithScan`: stream compaction using the `scan` - function. Map the input array to an array of 0s and 1s, scan it, and use - scatter to produce the output. You will need a **CPU** scatter implementation - for this (see slides or GPU Gems chapter for an explanation). - -These implementations should only be a few lines long. - - -## Part 2: Naive GPU Scan Algorithm - -In `stream_compaction/naive.cu`, implement `StreamCompaction::Naive::scan` - -This uses the "Naive" algorithm from GPU Gems 3, Section 39.2.1. We haven't yet -taught shared memory, and you **shouldn't use it yet**. Example 39-1 uses -shared memory, but is limited to operating on very small arrays! Instead, write -this using global memory only. As a result of this, you will have to do -`ilog2ceil(n)` separate kernel invocations. - -Beware of errors in Example 39-1 in the book; both the pseudocode and the CUDA -code in the online version of Chapter 39 are known to have a few small errors -(in superscripting, missing braces, bad indentation, etc.) - -Since the parallel scan algorithm operates on a binary tree structure, it works -best with arrays with power-of-two length. Make sure your implementation works -on non-power-of-two sized arrays (see `ilog2ceil`). This requires extra memory -- your intermediate array sizes will need to be rounded to the next power of -two. - - -## Part 3: Work-Efficient GPU Scan & Stream Compaction - -### 3.1. Scan - -In `stream_compaction/efficient.cu`, implement -`StreamCompaction::Efficient::scan` - -All of the text in Part 2 applies. - -* This uses the "Work-Efficient" algorithm from GPU Gems 3, Section 39.2.2. -* Beware of errors in Example 39-2. -* Test non-power-of-two sized arrays. - -### 3.2. Stream Compaction - -This stream compaction method will remove `0`s from an array of `int`s. - -In `stream_compaction/efficient.cu`, implement -`StreamCompaction::Efficient::compact` - -For compaction, you will also need to implement the scatter algorithm presented -in the slides and the GPU Gems chapter. - -In `stream_compaction/common.cu`, implement these for use in `compact`: - -* `StreamCompaction::Common::kernMapToBoolean` -* `StreamCompaction::Common::kernScatter` - - -## Part 4: Using Thrust's Implementation - -In `stream_compaction/thrust.cu`, implement: - -* `StreamCompaction::Thrust::scan` - -This should be a very short function which wraps a call to the Thrust library -function `thrust::exclusive_scan(first, last, result)`. - -To measure timing, be sure to exclude memory operations by passing -`exclusive_scan` a `thrust::device_vector` (which is already allocated on the -GPU). You can create a `thrust::device_vector` by creating a -`thrust::host_vector` from the given pointer, then casting it. - - -## Part 5: Radix Sort (Extra Credit) (+10) - -Add an additional module to the `stream_compaction` subproject. Implement radix -sort using one of your scan implementations. Add tests to check its correctness. - +* Tongbo Sui +* Tested on: Windows 10, i5-3320M @ 2.60GHz 8GB, NVS 5400M 2GB (Personal) ## Write-up From dd7771031a22efa02acdde2ba0c62ff8c31a920d Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Sat, 12 Sep 2015 01:42:18 -0400 Subject: [PATCH 5/8] performance --- README.md | 244 ++++++++++++++++++++++++++--- images/all-performance.jpg | Bin 0 -> 18519 bytes images/custom-only-performance.jpg | Bin 0 -> 19431 bytes images/no-cpu-performance.jpg | Bin 0 -> 16928 bytes src/main.cpp | 36 ++++- stream_compaction/cpu.cu | 2 + stream_compaction/efficient.cu | 23 ++- stream_compaction/naive.cu | 15 +- stream_compaction/thrust.cu | 13 ++ 9 files changed, 306 insertions(+), 27 deletions(-) create mode 100644 images/all-performance.jpg create mode 100644 images/custom-only-performance.jpg create mode 100644 images/no-cpu-performance.jpg diff --git a/README.md b/README.md index 3db0b8d..55c615e 100644 --- a/README.md +++ b/README.md @@ -6,44 +6,242 @@ CUDA Stream Compaction * Tongbo Sui * Tested on: Windows 10, i5-3320M @ 2.60GHz 8GB, NVS 5400M 2GB (Personal) -## Write-up +## Project feature -1. Update all of the TODOs at the top of this README. -2. Add a description of this project including a list of its features. -3. Add your performance analysis (see below). +* Add a description of this project including a list of its features -All extra credit features must be documented in your README, explaining its -value (with performance comparison, if applicable!) and showing an example how -it works. For radix sort, show how it is called and an example of its output. +* Radix sort -Always profile with Release mode builds and run without debugging. +``` +count = RadixSort::sort(SIZE, c, a, 8); +``` ### Questions * Roughly optimize the block sizes of each of your implementations for minimal run time on your GPU. - * (You shouldn't compare unoptimized implementations to each other!) + * Naive scan: + * 1:m - 1:256 + * 2:m/2 - 2:128 + * 4:m/4 - 4:64 !!! + * 8:m/8 - 8:32 + * 16:m/16 - 16:16 + * Work efficient scan: 2:m/2 0.1117824 !!! -* Compare all of these GPU Scan implementations (Naive, Work-Efficient, and - Thrust) to the serial CPU version of Scan. Plot a graph of the comparison - (with array size on the independent axis). - * You should use CUDA events for timing. Be sure **not** to include any - explicit memory operations in your performance measurements, for - comparability. - * To guess at what might be happening inside the Thrust implementation, take - a look at the Nsight timeline for its execution. +* Performance + * All performances +![](images/all-performance.jpg) + + * Exclude CPU +![](images/no-cpu-performance.jpg) + + * Naive vs. Work-efficient +![](images/custom-only-performance.jpg) * Write a brief explanation of the phenomena you see here. * Can you find the performance bottlenecks? Is it memory I/O? Computation? Is it different for each implementation? -* Paste the output of the test program into a triple-backtick block in your - README. - * If you add your own tests (e.g. for radix sort or to test additional corner - cases), be sure to mention it explicitly. +* Test program output + * The extra output `XXX scan: 0.0000000` is used for profiling the execution time of each implementation + * Radix sort test is also printed, where the test is carried out on the array generated in stream compaction test, and compared against standard C++ vector sort + * Performance comparision is also printed. All 4 implementations are ran on various array sizes from `2^4` to `2^24` + +``` +**************** +** SCAN TESTS ** +**************** + [ 38 19 38 37 5 47 15 35 0 12 3 0 42 ... 26 0 ] +==== cpu scan, power-of-two ==== + [ 0 38 57 95 132 137 184 199 234 234 246 249 249 ... 6203 6229 ] +==== cpu scan, non-power-of-two ==== + [ 0 38 57 95 132 137 184 199 234 234 246 249 249 ... 6146 6190 ] + passed +==== naive scan, power-of-two ==== +Naive scan: 0.042400 + passed +==== naive scan, non-power-of-two ==== +Naive scan: 0.041088 + passed +==== work-efficient scan, power-of-two ==== +Work-efficient scan: 0.127552 + passed +==== work-efficient scan, non-power-of-two ==== +Work-efficient scan: 0.126624 + passed +==== thrust scan, power-of-two ==== +Thrust scan: 3.615200 + passed +==== thrust scan, non-power-of-two ==== +Thrust scan: 0.016640 + passed + +***************************** +** STREAM COMPACTION TESTS ** +***************************** + [ 2 3 2 1 3 1 1 1 2 0 1 0 2 ... 0 0 ] +==== cpu compact without scan, power-of-two ==== + [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 2 1 ] + passed +==== cpu compact without scan, non-power-of-two ==== + [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 3 2 ] + passed +==== cpu compact with scan ==== + [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 2 1 ] + passed +==== work-efficient compact, power-of-two ==== +Work-efficient scan: 0.125120 + passed +==== work-efficient compact, non-power-of-two ==== +Work-efficient scan: 0.125248 + passed + +***************************** +** RADIX SORT (Single block/tile) TEST ** +***************************** +Thrust scan: 0.016576 +Thrust scan: 0.014912 +Thrust scan: 0.014976 +Thrust scan: 0.015008 +Thrust scan: 0.015008 +Thrust scan: 0.014912 +Thrust scan: 0.014912 +Thrust scan: 0.015008 +Radix sort: + [ 0 0 0 0 0 0 0 0 0 0 0 0 0 ... 3 3 ] +Std sort: + [ 0 0 0 0 0 0 0 0 0 0 0 0 0 ... 3 3 ] + passed + +***************************** +** SCAN PERFORMANCE ** +***************************** +==== Array size: 16 ==== +CPU scan: 0.000000 +Naive scan: 0.020928 +Work-efficient scan: 0.060864 +Thrust scan: 0.017600 + +==== Array size: 32 ==== +CPU scan: 0.000000 +Naive scan: 0.024800 +Work-efficient scan: 0.074144 +Thrust scan: 0.016416 + +==== Array size: 64 ==== +CPU scan: 0.000000 +Naive scan: 0.029664 +Work-efficient scan: 0.087968 +Thrust scan: 0.016736 + +==== Array size: 128 ==== +CPU scan: 0.000000 +Naive scan: 0.033600 +Work-efficient scan: 0.101184 +Thrust scan: 0.016480 + +==== Array size: 256 ==== +CPU scan: 0.000000 +Naive scan: 0.041376 +Work-efficient scan: 0.126336 +Thrust scan: 0.016608 + +==== Array size: 512 ==== +CPU scan: 0.000000 +Naive scan: 0.062336 +Work-efficient scan: 0.193728 +Thrust scan: 0.016672 + +==== Array size: 1024 ==== +CPU scan: 0.000000 +Naive scan: 0.110176 +Work-efficient scan: 0.346016 +Thrust scan: 0.017088 + +==== Array size: 2048 ==== +CPU scan: 0.000000 +Naive scan: 0.219840 +Work-efficient scan: 0.003296 +Thrust scan: 0.020192 + +==== Array size: 4096 ==== +CPU scan: 0.000000 +Naive scan: 0.001504 +Work-efficient scan: 0.003008 +Thrust scan: 0.025056 + +==== Array size: 8192 ==== +CPU scan: 0.000000 +Naive scan: 0.001504 +Work-efficient scan: 0.003040 +Thrust scan: 0.038944 + +==== Array size: 16384 ==== +CPU scan: 0.000000 +Naive scan: 0.001536 +Work-efficient scan: 0.003040 +Thrust scan: 0.067168 + +==== Array size: 32768 ==== +CPU scan: 0.000000 +Naive scan: 0.001600 +Work-efficient scan: 0.003072 +Thrust scan: 0.363552 + +==== Array size: 65536 ==== +CPU scan: 0.000000 +Naive scan: 0.001888 +Work-efficient scan: 0.003136 +Thrust scan: 0.449568 + +==== Array size: 131072 ==== +CPU scan: 0.000000 +Naive scan: 0.001536 +Work-efficient scan: 0.003008 +Thrust scan: 0.446400 + +==== Array size: 262144 ==== +CPU scan: 0.000000 +Naive scan: 0.001536 +Work-efficient scan: 0.003040 +Thrust scan: 0.649120 + +==== Array size: 524288 ==== +CPU scan: 0.000000 +Naive scan: 0.001696 +Work-efficient scan: 0.003008 +Thrust scan: 0.912192 + +==== Array size: 1048576 ==== +CPU scan: 15630.000000 +Naive scan: 0.001536 +Work-efficient scan: 0.003072 +Thrust scan: 1.351296 + +==== Array size: 2097152 ==== +CPU scan: 15627.000000 +Naive scan: 0.001536 +Work-efficient scan: 0.003168 +Thrust scan: 2.192576 + +==== Array size: 4194304 ==== +CPU scan: 0.000000 +Naive scan: 0.001792 +Work-efficient scan: 0.003072 +Thrust scan: 4.034496 + +==== Array size: 8388608 ==== +CPU scan: 31247.000000 +Naive scan: 0.001536 +Work-efficient scan: 0.003040 +Thrust scan: 7.794240 -These questions should help guide you in performance analysis on future -assignments, as well. +==== Array size: 16777216 ==== +CPU scan: 46879.000000 +Naive scan: 0.001504 +Work-efficient scan: 0.003104 +Thrust scan: 15.012032 +``` ## Submit diff --git a/images/all-performance.jpg b/images/all-performance.jpg new file mode 100644 index 0000000000000000000000000000000000000000..56af743dd6bca505012cc84a16e17f901938509f GIT binary patch literal 18519 zcmeHv2Urwawr-J8k|0P<(gevOAV_Es2@N7SCnZP{kPIR;pnxDjKtPZjBxjl&B!hrt zP;$;WNbY{6XXeh)b7tnm#dOkac+U)WXul5xheC+Z9d*bqg~F#YavST6PZBj+_iI zTN7^7I3NY!Vq@cAW8vc9;9Ri{M(7761`N$jgCk8rNrlk)k5r{FT(Ep8-J?cagiGIsF2bom-N1tk?T3oF|Vc76fD z+d{&3r0z+}$jZUw)zmdKwX}6~O-#)mn_E~~IXXGJxVpJNd-2lGKj2khP(ZS5VMp9ThphDS!fd>xyen_pO5T3%UQ+uhqg zIQ({Wd~*7ITxbBsZ^Hur{B2;rjEfi?7dj>;1}4t;aiO8Rf*S@gCKlsOY!XQooJaOo zuk-oflHLtZDQ>*PbW3%I%-Es-@-+zmEc5R7q5T-ye{EpC|E-bz$H4wJt}y@~0}V7D z1~C8yj!xN9y)XaS#*RRySZ0g~Om3{>U%ne-r}`9!76#FX1K3i7{x@#LqDqCd=xtMG zXKWF3vkzrzk~#xnn#MXsPkA#DfX%RbVgBaQ%^tSkSKYB*Xey&PnY0&tVparHdp1%E zL2HjoTp}R?ZPx zHUT*`0jdkjqF6^AB}mSlp>jpc#2fu_GQnv=wXe|S+U#o%^5nM5@HjkK&!BD9G+Hj# zppW$>G*yv z7YdMK2}1!s{qS9Nc$3+z-6ME%Vrj8P0$R`k@J%yXY*6i^g}0+3NX zdd`8fN+%e-DBuwX3b^2i1H=9=7@x_K7-20q7VbiXRojQzjw#n)3=Mb-y;CJk3aEga zhM1m`qMl6sps7m;_vV66f{c8CY88Eg{_hr*OuEGI_1kyexbTS8@{+a#G6THOusAq- zl{7%A^4q31bZ{`Vsfyzw5pzZ9jYund2hN-teTe?=R+>!n%#4JQlsh#y>ng7nF>QsD z+Ru0qzc?QL#FmU61a8?=maFd0nV&y@dQyPD2As9DLfoJ)Un0pI{<>Th+X)ybo!Q#l zzCXChjel)diD_brAG%H@4Pb@rBeRw>RV1i2lwgB{u>L%!FO%sC3T`u6k9-?Rxs?{x z#fRsTPBB(M>=KAU0pTlWT0!s?qVIq=VuY0a0fL)`H&8&bIEa%kbl@j+VF+?}bobJ$9Gc&y5QFR!E8h`shLMZp>>q z50f3MRYL(1o*=+tm!N<>F~Fu@JIQKS=>qKleoobe0y+*nb6cSYINzb%ng<0i^%`9wF+== zz@P^aQ-AZqKQN9QtvE!K-(*yNOQU(+n|5>)<_=1zH|<4xK1+7axW%_zvA)wN> zbBLhPeluw7Ml&!#sq=N&BSR2IVdv;n2?y|i^Q#<45^A24*=*1eOTds00quQ>Am)5$ z2(-J7BMN9+*9>wEFh2rUeX`6+&*`1}g~xJF8kYStQGjn(67sg@!=cSBuEhlD7WJ>%8pt9u3AUaArdQGiK2iL5Y2J9&p`rMr`Ppfik9I_OiElSCCSI`#uS3>RKL zk)Y9{ovC~?g3x$@*^5Zi)kRs_VugT^ngf^vb1@74)`8qx zw$D)jgo8-XJkr)(6lXksKXj5^fekii{c^`EBOlMLc3k0gmZoSz-t|g1AE>hTg`(4N zr_wCm>h+pIHLXA=g_?@$=!vnbLes48R!X*k8^)8$ZzzHfnDT$akM9^#EYN#`?`Do{ zQbGagF=Q3mam(laqPIX`WJLi33PHelp5m1|>xK|4FPVmrsM67kYk=(P>Byu*B2JZc z>D-M@EYsnHd4^@SzNW5IV)ckcn>XW2Fj@mr6;>98R#&x55h9pZl9I{{747sa2Lc}lLv zt*j*DLoesgt*YXsIWJy(FediSso8(8(txEn z@qy>~^sW81%@>(wN9zV+#%^jziG9#J2n9fAv>aWKUrlXGsQm;tP!H=*(Q7AvCAB$&98Zw_P&+afEA~O9H z>xOGxsh(>-wX(o-cjW*m%!EEKGI+~TYBC&9bdS_x%qbcUr+DOEKdGS}+ft$C%)b8N zi!avwYx)=Sqmfgu#K_Mcv{RnxYwsg_jsi|6=FLaP(gqKEI5eWQ%P#JYEzT{D+`zbU zD|mODA^2dfs8EzyY&vj@8#xrd8YweK5hHFqMRA_Q>(c%fE^fth`B?WzuC_G#5M=Uw z8;bD`%m@3W3m@dPn%XKu9`x750}}MSI407h%4pxWU;Q3*dq)teDM3p{qJUNARzPPL z)0c;EjqkC9>g3nw*;Im@*b?~ZR~$qAOP*6^wr9{fv~is=)`!TPQ1+fb(_4;k3_hR* zW*!0@N(L7jf|@V#dlnQ9FFzll4{Ecb{b?Q08@L2F8dP<@~$2*#J?#dNHKaxdPSpa5~r>7!M6Q`p8h{mEWA3TRM3k``)} zoPi*Yl{S5P)}=uodW2aJ zD6cysDLzs|#j!s7`>#RUI}l!ZZf0qpZ4n(V$t0UI&ehZtUe2}BOJR593k!2T%Fa7~ zl-6r!$2`1P_kN=0gjkUce+6PH`;EX`c9hIJ5oc)gfJ>L!x<+G2aSWF6Nde2BBFTyE z-26BlhjO+uoBNy+j^&?`VF@?3rj=EEtz^j!G4pu(%H&g)Dviw(Z~61D%ki%v6GP_) z8LLN@rZL3&hVQeu7Ais|>!E!26kDF5&y3XNFCq?P^mHODnoI1Y!mA$O_7-D$k#($z zUrIy)4~977w7(fz3m$!L($eD1jG(Z^VBDo_6RV;{@Y@mT>x^kPyL$*ap@4~xTpvtB zwcnom-vbE7pW@Qkx9bDA2TCMscM9H{=k%tvTV_m*;}wb(Vvq)Z*pc5nXn)o^D9$== zmCey;emPGQ{h~W2r^du)9IfM{UCh!lUU+Aw8?@=RIGrN9a4pB|tIRT=M=@%MN|xDI zig`2C&ce2iF{5YvRfcR1_q$WSl4(J}yW{pkuuqt(aJ#-1 zljkg#!+g_GB3;-S&&HBShX0zZOL%ui(87}wqR+0pT=7oJf$2NW+us(u8>-?y&#S#w z=3zcqQ?ryAUnn4cN6|LIDPl@n$v`z05GpE``PFEG_Q8~Si7zbOmy#@9+-OSq)=_Dj zSgVL;oJfLHQ)$!eT%@wO3@rtnj_YuH#3?qzNq7uo|+QAsA4L&b_oKO&xze)>Inx; z6(SwFuE>VRmtLcFP*PKNdmyIpsXDx1lfY-l?%PMMck)DI8SAXEFFwP);kqlF8#99& zb7hFIE??%c*AY*}Gvm5`Mj8ofr@x4ozlxau&YmpL8icsdxhUY_IL`&OGbqJ=J^}t3 z4UtK{sf*zc5u5VCMAx^pW#@?f;Z)#1!QG5GKl;w|EFDom`;A-z?UB?H_$)0b2og_| zkObWD<3;+N(68`bPms@G$3V{=64EBTwjFh2IZl*oeJH^Scby;Dx4G?|qVL~yaYr~`I$G^w|N73XYq-JD{@@as*bnte7FYo9FlmN< zBY|H)D)XxsYpoqdEYl`fo_B1su%Wf-ZMX_iZhpRz;ra;nRj}*y&V*tL%-#l~qHoZ} zxe6(cCB=3uio$|X3x=lV9vg5+9JN?`@ZZuWfVpqzCOw0 z+Ek-{7-&L4F)KkoP;vkT>FRrA6AYxQ(FaT0EPcuOq_Gv!r0Q?~o;Q;HB7@4;uQny^ z;UpoOv{1mQmH;Ru|HxB+<*oneu_o726evk-da^C!{4%N%r`RA0?*-pq=*fNbEn$I< z#2>D3SOozsXFQ;azWH56=NKX7(l+Wr^_AQf2GX}|_{sHhkl8Ag9A_e5P(hKGGe8v` zB;EzeaQyEnM0QYyM{wBw-~Et|k+KBCZ%0xFN~*y4k z>yo_Ou<7H(PvRrr_PXIy#yJfiKq*p?)aupdEKcnzV0~2z1-xkwY}kM#oxA*# z&-Df&^y;`jck2jtZrQ#RY0aJ3Fvx2;vE09n7uFW3*2m*4SXmLNLI1Vlsw~IHzQ^}2 zt4G0Z1Ef(nQ4q5tcZwQkv+VBTGO9Q!e#cmBnNih`7eJb60A;AM2y-l6$i8&g=-iLrbl zL321`lE{YM`aM1KtA+*D=D%ab{(HFk*YoD(x50d7l%-|#*jo-rT*A5QEm*KD=5?Xe_9IE=K;eaGF>Ihf z+xoO_so`nynv0-IP(?AP(iJ=A8B=ta_8YH0|+66HCA$D@5w^UD zjLOv#i`cv6zvuPn*WeqeDzB(5kCcN;2CRl%)^B|Hnk$NI+)}{PmI)L zBV2DL)V#NvK?fLUb|kkAHVenZ+i}Lu3(HQdm)(WGHqZ-B><2GX7K$M|d<1NC4xw1I zx;l1{aGvzUuW|QS7bzn2vsC;r`XxU?f)4Po?hWRvirAI36(#GzV~(US@b)5vRvVc@ z|IZ3rGo5oJ^Dv}$!hec*Gh`D5l%fFk;kA-5uBSb5E{Sb;D?Ex?me!jih1Q!Vj=Op)LmG$_9D_0 zA$aWG@>u}eytfcp2w~|Qfs9_}Ab-BlSfe9zN!{>%T2)9;@mMi02E%;(eBxo>$H#)Z z)GzPFC+}YqTpHp^t8qYY8UVhEVT&m}n#vdw;NW5( znzSNU9+$vXgXEDpJj0KhM|c`Z)r}N5VKofqnUJe=WhwBlL?~Y0#3xqSA z*9`L1J#YFlyn_yYcF71#HR%r}lGRr=45UZT&2P%Z@U`1Hsq>6oC-8&6MeDuqjn*cs zn>{Nc+@hr=t5#$2Ak7OaR=XBtit=wwY2RO**dpVVag_d=*7`T@u{C}6b@UZL4P|Fb zEuJj@2`WvNt7AxYQhWOm&-u4k;412QK0Si4T6_NFbIiv&^j~}>-|=cMs0t<&Y&UqB zU)?Z0qV}PGC*U{h1}U%GrQw_Gaaqsie*BKX1qDb~h;C>Y-g$~6|EJ;YyXS7RX{gO* zS{-ix%EV05*5ew3$DW5WPNGtZ(b`yjTJ}o$73SiaaF+c;kB*FLw=ZdLm^`!iGPD?W zZ~%LQsQ_oJEuTySJT0QyT)*Ol4u6x{MliZs28{VrHqVo}+ zo$ChM0ZYE~&Cm3|70VLTMu+pO$iUQ9z5OVG8A4(7s(=B;QIiqd!1sRn!o(fS=kfM17aopC0 z#SPO~!M~0a-;7wZFS{L0x~4v)N~P1xd}(7{LradPEX%f#@^*JF99bcq>W*?fEgulbQ&AA2zplU-0y}z`C%a&c*R<_)RL^An<&po)nnJyh86Q}= zkgjN>$cbzdAzGK9CHs@2A)U3el{1`VO4c&dT{Fgbu0ul=M%l

inz8dVe~Hx3P8VCVAkQX`Q`6Bt%rCB2&5g^Y z%Lf@1K&M8g(Tn4SeXH+T)Jq= zr)yvIKCz&qQ4%(Pm|6MgGh{!CQj+E03O19~~txjCeloLt&TMkpq?Q`I&M3 zgRX(UDb2r@-v2dw4)*QkUG5nAqeqqR_x46nfO)>|Cwk~>YGkZ-*tusf3Lw@Z7~gyB zEU3^!Ri+{>nf#xG=O2@JWRY_Q)=1(>^kwxk=$L54i>~QXkp!&*qNCD5gtew@;DOW} zlyK7W?!q|bH{*xzzePQr3VJsXE`Y&t=uN#t2+gk$`f zMi`+I!Kgp~`_~@M<^FnXhC*Z|*TK-msmG!BbF`i%BF(>OUD?e!IU$ljY~Fdn?dA4oAJt6V;t9x3_a{fEXU-vyE~n>3Z0s* zZV;DxH%vi!gl3*C2Ps`$Zpi1qpesjSo{E?C|(R@HbH{e%Yp#-;@IuT38X7Ne0? zRQg(K0+D60jCq;2?lp^&Ibg%t+jkQ4Ka^?dTCH94zYw6_8%Z_WmpD_;WWJU7k}hEd zBlC|G^UoVs{;B&)xDEIAbs|2G9Bvx>Vn)cw4Zprs3Wz_n_kM}Io^bI#8{X%#j64U? zwqE`KVaK9|kuoHX&J?+}4+9OmliLVI2F6fVvth>r193is8tXfrOHWgL&w z)=zI~{-{d*tjIMAO>)0&8)v*};z=F-W=Gr8h>d7U zj5Ons2^!+Y1>$369O=J9ORxWfu8i838H?Olh>&Pm@DJ=W3DcyDS2dE#>H6lL z9H8X>u<9{!caMJV7wtjjhQ4e45W0riTtC) zkMFl{;}1%sxAvTTrzFRl=6f2^`5u$R!N3W$6-k>3nDX$P!MjXSh zcEkGW#rt~=NkV)13(qe|N|`s;t*Cz zAM0R7)CVYl_kU=voVchmS51<%Tcim>0zY zzXoel9q0hyV^N(rpQ;k5x|L++Q6 zA5#UgaQu7)*GC3fYnC94l2rR2wPSCl+oiU$BD*ya7BNjRVE&PK zL=`@In5803hW%1YaMjfYLc*UD%6nh6?1^}AI`Q%M1V*~c?Rut)zz?R1u&0k>7rkt3 zdN~)0*C)=j{*ZzEN72Rx2q6zMUwQwm9_Kl-3CA4@gFA|!N$QqVY1&_;%HbhAWLG7 zI^^W^Pl|?8{U=Et<7SEGrAW9WjE@Ubo&+SfO~2;cM&RjDD==$LJK=ZOj*UA#qL_bX zR{w3vdpE8I8gS4RA~CIC+h)AVfp+Vqt5S8h_ogu~uNBH4^jk9GTq6Pec zK_aYg{~KyzzgrE_qH-SFsL3JLM3)lU&)5R!kM|itDjV#MWRE@5{hC9R7py?^iU0f} zSF_+5BzQL0t{VjJA0=?K5}A^g-O5j3!F7o0^Ugu&DOMZ^r?7=86kvHt>NGq5&Zqv5 z{_7fLo3vlUy&~A)7q|%B5dz9C*oSLS=qmxsJDXRUlTxGGk@qgZ&@Hn|oY#Dhd}dhG zdzKaW9m4a;#CdA=XrTey=S?8X;S5@|r5ck)M6+cwEx4QTv5Buy!p>VCx&6apW$G`@ zo&UJ3sQEKej{h9VLug<{^S8#Nf61PVW}msPwzM-LJjJCwnUWv7%_ZdQtFuUbz4LsQ zc)t%Fl;5s(Zsl`<-74!FS$xY2@lqHR^dxvLEa)Ro)an&{BmSFhn^Wx@?)N-Aj5jMR!ux9)o-~sSi7* zOrrn5p;>Y+nC@n+$jRGl8#=8?9kynz8k9tRNXz5%?0al#|I-5IuhN+;Jm1Km(LPiA zOPT$I(f!V4x%VUd>gFaDOsw8GV*BGbhNaFndH@Uy8IYXybi8s;xgZnaAJ=@xh^y~ac1MstSK_2LfLoPi4Bz2#u@vz-zJ_XzFUtdfCSGAhMOE_e*<=SESz$MRdy z6%9~CZg)e^U8?DC;`@uO*Wz#X3LUoId0o)x{AC$B%TWowGNT11A{%(YwurGmoLE^& zT`w|l5&q?A*=tkG!`x2REBq*cV7+gu)9$+7r~Z5Ba_WA@_X1KHSF8&^)Z|$hd|?mG z2nBw7#P-i$DD3CMY6dAl5|d7nw9LQ-@1L27&+sNKaJ{%&04{-Ek5Iro;sKOo?RDP7 zXC(D#NaT1^xb&iQRJiWl&9GaQ7VSo}=clJ8iX?3i%PWj+6h)1_4=k!T$PGf_NRpbc zQ!-~1aJCOJJ=(NijtfGAxNERRP2$0R%z}BOp8*_s-VC;-eDBZv{jqgg>W|4?5CFeM z8nPxbAgf5H1ORReMmnp60YJRF8(qf9gAGT+CO>Qz)WjOW$ z1xAGUZxbT+t0Vkf&-lZHl>K3b`Uh@dsdh^!^T9;Xq(>iCZ#W2gF%Ms_9->=res&P9 z@RaWBJr#OxD7yeV-2LLvaeCW}%<@+H_H`6+v_y{-Q9no`S&JSM2&M|~KU^&| zsR-d_)=p(yxQxNeB002hKsI^{i_#eAph;YM)8Q*g=}ahdcRBL-w`U#w@+^dn*OBr- z!IQb2`s$E76rC2S7?!iuu+_K?woFRM)|-qj$e^sCsK!%l96y0`_$UR~;l%UvGYn`M z8PH(>Vfw%i#9(3#be0*h$y6d(Z|dY)f*J3ncFFZn%CyIR;e1iJ>!tO_kjCt%c;3(M zuKoFXAWVu?!>W^ZkV)o^}&>|`%Zy*G;JQP|0< z+8vWqM7qA~0H&nvwD9Hor;VtTpYK z%eM=9&iuR0UEJF9IRX)1oJldl!v})y{=J4>zFH|y*GK84UmE!f&`va}f4tGk@ox8}jVnu_8O>7|XT@0*xBo+%;wr}WUHcCD zDd1Z{VscGM*$~4Xk4;Kd0Ap=c8etVJA<;9Sc}HTejMC;Ed#9bdJ?U%u(+sG^xRQtYnT4TG^nDQKWyrXlUyBn<6>FM_o zXOT4vr=HhyJ+gn(#%{DJy87)+B`Zv@qF`5&LmJ~k;u4+X@ zWLd6eeXZI*#3H)dAfx7E4De{Xi|izDtwV6x(V=C*zH=Q?RvxFuS-vakFZkM#f%Q?F zmUsXL(^$5So1@gAt#prnNT%**G<0n)6$*?YQ~W;cgzwtIq2bDg-?Z-`tNXIF|K!5n zyk{UO*bW)hhVBq?#2UaUy) zpZ+lWlUIhU(k{UzZ)okKjwm~N+{;FK%~RK1ly64g5S7NiB`_4+7!{3w5|FJbhiU&t zZbmC(*rmACt>C(~9F<6fuG39lW(c4_G&h6R`;ZL-uc?E}&N)GJd`y;;9M?4BlVA`f zo=mziUpKLLn}q0&$EbqF#)b|5vFK|KZJU%%J8p_5ZY4FJsk|Q?7l&h>gYm`y_n12h zaL)ef_fnjxGuEMnFq6jdjo7Uo$J6fnr}c$7uku6c3bHLq zLEW^{6ftjK&J|VpL^eEh$G`bT2a{Oi0#zY*_=jGvn1R)A5r!<6C-~V-^_`kcL{0m! zK7Fr|b?xE~J_xxn+$kvE@i{(;*B?vGVzJ~4wZ}5e5`rD^Ey*l<{2~$gd1?z?%3ZAm zq8}Q7Y3wcZs_ac%V!Yjx)Uv@tE?5SY?eJb-ehnrePFH0Vx5PYyGZQPk36Ra4mb#*1 zO_*j+TVzyS93V>caNL`=ICS_K523|V-A-%5kj!>F4(tdBLRyw7g#n+0{i3c`FnZO_ z2LEX@#aMV-iiT`5hK~fvcM+vUnQIw+=R<`+VuDpz7WF&x+97?@F|4t>X>?LiDt#74WLA_|3&0AwKpY?`{DSPX{Ufb2^+x3h-2d~xd71u?0)X4?c z4qRzW(Ohfh+IK;_8!FrNbwpj9vX5)~+7>O5V(6mYv+n5d0D?6N{N*XcZ=rIvrQxf! zjWtWA?WhMU_Ef*jyl93TUeV43iTjL+eH)+9#pAIf#hvuid*5P+$CU_rCgwh|xTQ{L zXH=Jjl3tA##YYF4?WfnM9=y-J&XSQiW_X<@b)u;ROFZqVP)7Mc**XVH38Dz$+rM9p zo1zg?I?9eMoQ%DC6)%MDvj*%6KgV8&pSJx3)sHY%$5Q46JJWtcl^^eL{CA|1e$#P| z)cnGgw8Gmjp;DH>|swaeRT!zjS!_M{=rfu%Z_-u zC|f~w)Yp;WhO&Fu_x8NLCh>;Gwrv!+I(Wu?$W}2Kx&3B{}WCXw2&@^L((bu$v;t!V#j2 zuXoS%u1-iQ%Yfv^W?$j?h6k1vr(Wd(0{0O0){JJP_Fy9-8`SN#ElCyQx3C!MTcK_} zfb+G}I=`WjT}h`1v7UypN8=q4oIaIwFQ0dv1Vs;=j-7s)(7Ru9mv+!?HOe^uN4n5RDyl^z?Ye3$R;euOmz(d>^?|fq+h{}fy;hu z(_3?g>H?qP|2aT3Zk_%&m8}N=*rlCy3|Cc@GQfM3Z77`sTWbC0 zap2EbEpRsy{bR7rT1p9lt=oq-x%10Acl~M{zAAgOElum)ib9@%h_z1 ztDiMg|GBh_po|3OCyeE-+-||{E+%rAOYX%uke6>`%UnZ1`=zJShf4FEiLi!*0{#^T zEXN9}kk*D4k=PBVPG{yw=?aGNPO9+O$LwA%DrH9u%bD0|gJoiB&G{*lubHw05WyoA zSM=n3(ky9BKiRS4V#y!kP>2tq4eGfCz-kRfj{5JMV9zs&%3RW99mKpMt1^{vtu8|m zE4|Z$IB1JJkz5!pO&L#6Ndk@X^PSPFuCro~Jp?L)CXK6ycx+Yoh}oCG7a(pgmij*` zet_-HO1XXgm`62nyv2o^$rrmkZV+bkF^pTwQcq|B|E*H0E9l(Qt#QoVS|sfrFMYAA z93eb(#)btoZ|zjps>oawhw~7xT@#y3IOUf3&L|}rFxn(K8fk)Gtm|`42Oz{AeAsu(lgJHL>( zoImlI2nnVmqP;_Bf$pDDX#LDVFwNe?nl8TL^^>Ao6=Sqik`!-<_ZA{6x$n64Sg(%7 z$54I9$kb=&YE8?{zVE$L$xNMif}c6**k-TagaYg)kHZTQM)*?`Qg*`Eh?#pDqAnlo z%ki0|GqHy6z3{%(`K3tBalXfJ7iY_6V{@bHSqBvi_x##DFx#Jn@h6>Z|AI2(pLgw| z*NzCfk}Df1Sz}gQJJ~aHYqZ!OwG<7WM`-`9-cpDal98J`KGt}ZQvugB_y%!-OCfy| y-_twzamBy%(5qedXT!ivlWV@|ayi*7u#VoO8}mkQ^m9O3pdw z*!`A!pMCGqv-i2@+;_(t{}}&mN3lk)RaV!mnl-;Uzd0B38*&!7{ym%1qB5e*zQP*@=Ndt2=f0J1O*!#8xI$c5+9$E{|4<1{(t*}Yyya|P|{KP&`_8F zR3a2KA{1mRzyJz}f%4l2_}d2s6%8HZ8YUJt4ldZC;yQqef`*2Qj)sAOjt=(r0-pou zL>R<3Zp&OFQG1HX*Dm(>?`EJPSiaghfQ} zh~2v{D<`j@_&`HbOIt@*4{Bm+_UyU2g{9LAXBSsDcaPU^-uk_J?;j8m`5`Jg=HsW> z)U@=B%&hF3+|si0ipr|$n%d@;*0%PJ&aUpE;gM1J*thTFbMp&}OUo;(YwLUa2Zu+; zC#PrUKjcCI(0&&S{QJ9LzsN-d%7u!Kj)so;LoO6lSFoWGp<~>*eT`T~4fCl32^0S- zEYiE-DJ5U9nFZ8$$&4KbamZN&=UDfCi1wRg|C(UF|1HV>L$LppYaGBsLjgArjR=4M zC+F;`J~;oMYlSbh$9A}Rc8p6(0+s|X%xY@+AF77AUiaS*)5lU_A&}2$!n`M|R@3lt zP0=-T0=I@0%lPgbsUPz)TQRUWOY*8o=&QjgT*-!e$?fo{+E7~8-?Wa?vZLn!7h8A* zFeM71B$B2JCc7tBAE!DrU;8OPA-xOV>`v@Ws@c3Sfita_p=53@L-R^66JSby=HK!< zSQjn3e7Gf#6h*if$5d`yZ2!X%KuF>xaj7^=y$Es?l|1Q^|s zz$PmU31m2_UYjCSlg0oW>EiH^^0Ntbo-eaivB1Y!e`;a2-TikfI)B*ytUQagJoB+} z>yW{WsroR1>@L-H9p;0)52`4#q`x=aXuOEi7O1_q+t${Tr<)RE452%Oi=4JKeI!ZO z*ENtp|5x{;{Y5wMm!IDeV$ZlTw|R7{+2+tF_loDm3pH7DV^TF)Kx5?aaAZpz(?vSw zx=NUM8%HPRyaq$CL23ZN{>-r@!*l~sW4AY3qm%iWdV%f z14Q;pmYOt;mWtx=u;O6;i*HluN=k0CI!}GSkn*T3X-bXGC!M3Mf<{>=1_^|(Ug!kC zRtbMta<~yf{x|Es^%z{~WGN5;UhBcm=tC>qvJgJ!5Cpy%5?J*3VYUB@L2&mi)^3*B z)$aoy+TRbRAvtxwcaRt5e_RCa>i`1?0blYPFGJ%p~^_5zXPS}79PmjrAE zb(1XjRIX4CVVBfBNTBo3E3XZ5i1`B$tayOTVR`eR?P&$75z}mE z&hWEQ!k=gQ>nwku($C_i{UYjbX!G+R#_Qv4pEI&1ETLvWYL9ydxT42)I61KrA8DmL zx=ls*G60y)S0+T4~tF}tQ%MSvXFpp zPZB~zdy9AObMQ22wQ3e-e^ms@DW+OdbAhTmW%{H|D&q582Io2wAUtIwOmRClG-`=8 z+pJWy%T{r)B5Nx#zRPlYiz#tk(^v8folF6~e*3xe+ky#BC7L$c_FIMAX4e%4Usqp4 zTW%+OR4v-4Lyzijn@2#7^Bm2~e&W8|KKo_tf>^e4+_wi0 zHU(_^r}w{pP2eEJce~>WaVBk7!n_W>iWRWS%I6pDz8sQLWWy1qk-5Y2q6t&uKa5}O z%B7!#01|l02HIMsfQfu%Cs3oQK#h(>0&A>oz^^)8xj}Ldw9g&>jhie<=ZbdHG@uOx zewe_2Xxe1`{&SU|V|hy^6|L+fTuvygs!Y|TCQ${dUp4ScH%)Ciq~sViys9{Gv zoJMLU%3|LwC&N_Y$(CnPKn9PkeVfv#EO%Lv{mFAcYV=mu!<7`wEDvDqpPj|8@Xs%G z8IeHZ7+=vW?d>bl!EA?|HSP`jwX;@G8_m)i>7V6gv|N3scRJMMyqr{DuUE{1oY+wu zX&Y_m_T(KR0fAI%_QExUh*3(KEU9*fsG2M(@XwkyYtdz>SkcvM>&8QyQ5*PAUjg$e zOKozlJlB^tw!QqR7FyW!--R$QS*UrhFg78(tdRC{^BhY=I2JL5xp< z@I7WYB~Dt9|G8&|#fg#YCDC-!Z^uWVagK2iw#xeybn|Z^gxHSX=tS!fTj-FFIWf__ zV#TZ5Yht0llUJ_sz|iRE#J^hqJ_5uHOa_~U=XE28Sg92%rIzG_N=`RSYB4|2r5e2a znDEHG3mVbU;X;_%;q1v0>Om7@&gjnUJC*Q8U7_iEO87n6i$qKz>|LC~xA|*(mUgsb zW4`Ot3ti~r81{R~RA;_nFN3y)xOaS~=02AVAEB2Qk8mxN*7_G*wtvet>9wiGU?sSN zCyGhubf0jv-pcB^BOLC}QrZHuvv6Vw7i(kq#9}G`DY(wJlZLAB^`jB)BPEvl3Lf}A z$x_`7!&zEes`adr?K`y8D?6SfHeOF%O#}xkD!%66U$pMps+XL(zxQ=Vmk+u)6kT49 zo3m09-I5)oi1=2be4hwg6imX$jY=GblP-Y~N-$*s88_NE6gOn<+n3sj zV7*!vaY_=|UI}u(GmE;Mv0x}R-hQsjVsEqXDZ;%VW1lWq*DFYWf0mx(ur znMV2oB_^UxtK8oZJXO@EG2u5cH>q*!*G9`hEwQa9^KwX@T2(B_hZ>Z84VSC@!zB zEx6WwU#Ou>)Ww&nHJ2RSSAOKE{n0JYpPgF}QQb5U(gz^KK0pFU;K4HJnf>9;{>fE9 z@xb#z8o~`q5~^~$14Pygcy;+JBv9vScix?LK4^B)XQCN%5lqSuVTT09KUmdU5Lcbh ziy(o|hAM~oMqk+0U+~V$cz8p0mxD>w^l9h*z=prN>cvbm6p!#qeZO57GTUB4QP~iS z-<=*t7`+isZvA`&oj3Cz-=4R-z-B{4=)kFKr~k5UIqQwKMCIcxc6(E^zW%sX6FX2_11|dvA3^}z@bs%_&n@v!&!jP zo;Dj2@mi_K*Hb(@eAFi0R`NsA+M|KN$a|$zC*x!;NWk4sWnEb*Co3-FvHaC49ZX0l zLMOwH*=6f70+wTXU79&1Bb&>puedVuVvy7OiRMwNtIlo9uU}+IHLm|Ng#1^YY5(kT z|Eu1K9yICwi>;eZ_`^LlCI=a}o40R-ve-E$ZyCJqI+XJB5i-xMIu9<4GaYkPCwot( zR-r`xjFBDmw^@9Eqx_Y<#~-8Z?3-5woeJvoucGhqh5JeG?^Q24HMO&&{X@K}9rbqK zl049fR)>g_Sa|a8plC;KF}k+hzv%AgUpaJ}>rA-0ybXsNje&H z;2q(_WZu_M)0lhq;j^H*Tp+Dw;E5(6<@-)`wbO%-DM3~*q4x?q@D!$OGBVzgR5%DF zOGw^J)#p7yQMp)pkBCZBQ<;^Eg8B1mgASbD$A4wf|5=}Ze?v5G5|R^w4ib!{|fu8tj_^%StFxvayt7ij@LEH=20c<#bc ziezR=QE2hhB+SBMufrO?EBGoNn>6h@UwmzR2^#vrmjXgOxgyJ|XlO_Z>zkAh810%37M8BxRCF~QhyMV?&I)! zCZI)MpEJ+AuqXw-ofUDa;wm8=L!KYH#LqJ_&a&Sm?rjhr9|^A;XGbG!&U35P^j)2D zhgrmIHd=A7V2OCfop^rVIH_qD=+oRyjLgi}fkvzy`(7QrjMPG4 zr=$0em3oH~x6UB@g4wA7Wdh>&=bP_u$LQdK9Dor+CZ}9kd1%ZQ7;ZW4tBm0zfimr2 zS;$*zt$UTK%eo+? zvPxOeMHcTVk&~6}BVg81@zN^wm%tk227d9Tbx;bGGQP}LMKO-;Ec-gw#Ev)jD&x7* zE8KTs^*(M+2fob`OTIim=-E<-54Vqxh)D5--er;QLB)c3NR(W3@L!`r0`Jf)#seP7 zT=Nkt!j-4>rmxyhZ47QjO|8kgci30axmIm@2G+I1y@oU0rf+qRXi zi~f(HvnXzBO^H3xY)=-s;NE3H#L#bRv*a1kC&Iz$&`_2*={kT`43UspG9*2 zqSj>zHgo&*Av^qbjUt_#3>Ez>oIt|R z;zsbb(ClX5J6EZq7272iM{VSAc-1OrN}(7MFc@WM+xs^0SbqX~&n>$2-qO2h!{aUX zun2EWbPRytL6u}Snv+r${Egp&&gZIUx5}a4+bUV_tm~EbzI?t!n-tSv377YCzExp^ z!#mqTQ#ufS{Langj*(%ynUmkK>@s){+5XQ@@Qq%i| z(fpAl`EwY*rp)mLJ)&w#mdo4cS+04BpcrcaWajz}yOZ}_S5C@fy1@i{nsezvt?4e@ z`=KMFN2&p2a-?{GxTYN>=ID&tMr9OXv^zg2OZM#93kdJ~9n6!hr;tcR1XDVkC$>kn zp+@B>ylM0GXt6y+Li7339vJf4_9B6WLda=us9|r39+dFRz20Bp2*<>{z0bWlP6xBO)6IM27c9fl*xTfC07AE*wtEnvl<=5-a5@6>Vh|)0g@y8< zTCQpt;$F&siY$0XGa?eJi66NtyzNqFkJIb>x`W*^QMgT1FC?m16Ldrb1B=mC(XHHc& z!g|BK66aQ`q}1O4?Kv#3apNoP`E~bvuz}5Tb35+CKy%*I(tvrPZzd(^cLT6-uG33T z1T^T4)2g37&=1`FS)vq}Db(FmGg`>2&7%*F!q0n8-J$onO3uOC_vIa?KIV{@mP zTy$Y>yNHEk=saDt7iM@cxySJS(1vGveA~4=c8Lc{V3=eD}c@Ff=E#&AO?$R)#kq1uoy!E^zXr};@OA} zq;XF&NcYTUYbicSj#UkI_lWzXCvVMWeV=Xl%@0nJ?%hsBpY}@jJ+j=S6CL-Cg zbIl@AXBj@Z|5b^;p}*pV*3b6vAA>jMMAGW22r44_JJ-`#0&8@Me6hAA^S;97$oGWW zrh2AMVP8mq&Tv6^`W(3sz(XM}hjms%he=uVpwPp{1X?q?vss=X+^(>x#~h(w*Z#h& z6tg+EeG+Ji3~^eMrjWRkWncj%C3jVNR>rK*8?E|2Mfqzl87WE=I^F^sl|<}l?y<7l zu|vhBs;r|^17ld{>&T$nT5LfsWX8Ye%v*P}E?5sn3b~QnURVlWzMpQB};tRIyv`6=wt< z)t|c%W;8`pr*X=W7sn?=-X+CgLPBK)S`n4!85CuP(=-xf{Dl{L9T?r7vF1J!3 zeyK*+&^HRtaHja^f{Epi>qCl($(Cc%q5g6$Vcu7f*+0Wsy@xssH%tswfl6u-XF*7m z{6af+o`KNn=7y<-6HijEtgYfWCCjVl2Bq)X70!uoOVdnKe7ZSs0~@cHGCi7#IeIRO zlS!d-q1RZ?Z%LG`QtU-RK+o=Kg(1$gxHxZ#43S$~O*qT#D93!N@Q>u3Ec=fX`|lLo z48NZ3s>NlW9B+D|ik6-Fai~4pFkL!%w!~Io$cI5U@s~`V86AsZXLrwf7E?+6&#cFP zru2SV#AjWI$f6bo&GPr%uo=%>!~+JzIYUU)v(=f;zE%tHZ4IWRW6bWP(^bqpA>z?= z3w)B89uTiyIXwtq&lC7j0mA30<9MUE?0sr=;yHQ?!I!F}L3}$MfqAF0&ldc@DmVGrXZo+hesnXu3|&;LFPIJPePb>93cTCUmKHb# z`4FF}2z6WULn&3P@ocwFNnHdTX4)95Q&vmC>1%D>t-=HI06(Lxm+y&I*z>Q;(>&F& z;?bw-lAu=x%-zJh=qW2haIx~qZ{2)bVEmzz=|xnO-#m&*HEoi{Oo6eR?{xQu>Yn)e zEt-nLNSX=39q`V;vXFEoaqc`BZ0TnV+f)Nw=FK&Ap81ap<*wDfj=|nEm0 zBST*(Bs5wx(2D*LgCHDyQBhqM45_d!p*wu${E;Pc&Lu)5F}Mpsf3&KIe!GVyhnOoO z%+y&GH;j+} zBG&aCxl;G%!txxu>lY;3RT%dztKKY()aT_;Zx4KOY(od8z=qa=R-P}-VOw`g}sJuT{YEq$3dV(tf0{*sGuWpZ!i}ug*!*J?I zt?(3y@tD}P=zX+bq#L^VQaa*%Z+-SOX#rGwkj@hu&JHOkQ(PlsNIJT5{*TCPg(q^$ zq;=>BbrPC`1YDR_7=p5r)IvrF?jlTyx~`Tm56FVk*lTtdtI!ZeL1Sn^_M2_6y`i+- z0{k17ebjCJ+EqcI`abVxJ^&FLED-#-s$jH@64 z$s6*QBh;mGa)b}McfUh+gV$C)+HhLE^Y%fava#eq-~z_GaUf)#8t_{BFEW>ZmC^kB z`>32z)m9edaNX4S(D0?8&g6@1GqwY=w{tIQm1c+vHV69bg;-};Zc~OKfz*iw9KH7o zbp3~g0spmKg%u`$-iu%@NR!T-pFeC;ztojvXzzqG_>2<(3`I zj+Wb^spsvb((Xqppt>(t=CQR1XZEy2tGa7DaAaz6eXN&JWVP?>p(TojxGkO(SpHTW z@ZZP;_cm5XESqHN)sXNnTS#}lB3!FJX7HbN;*a6rAaZmexqo%3Blc>;T|ms#Y9`Vyqn)@6-F>;Y7UWl)Whyu^ed*`faJGWk(bK+%IU(z+x_g zV!rntHI!;+AmJ2M=zX{b2TL$Q1HT;{hVmnUEeUU-7i0B>vhwEQ6!}}xU)!fg0c<=g zY9BtnYmf}B0nyznUyxKFQ^_0*U4sfGq&Q@7fZz2r_QT&D=Ev0!A6?<67AR*-d1r#r zpy#Wn5=IBz`%5rFq4!r<1<>vxDWxf6-tFmawC}0Gjh}?_odb==j&k0V$HK;rQu6+v zgEEFB^x}e7{_}`LqsV6EzDk*kh>iNUe(U+t!#=qon+ILE=$|sTR+fJ6p$G)LZGzlM)=K0)ut>Z*#bRLHoaVs$6ZEPJmwl3;i1}6V3zjj!{JTk zREbE%6~l+F$~=>{l0I}FTo^_obvTUUJbP;tS$pUj`BDQodYgYl!KIm3U~Qp|M(zZ( zi}R&GX>L-NX#uoru*}uLJ&)u=tyu`+`TQpw7_s5!Adb0IyYrE+9u*85Dr?lE326=( zLok7xgpv&W=}hJ`>z#sXAo=wbG&cy`PiPEAs$qKBC(b65Ne zsv<)ZB_C@aG(DW78{RHs6xR*HBgUTfz%x)%FbOdCF^&`v@N=<_-Uv#>ES_~q#-S$< zU)uExWC`hK{wCeq>!|y2j3xH#ZhUM^lrs$!ii?sxvlVYp@*$~02@c^!i$?-fh^T>a zXp4y%Mnb8zFBY7Am}N_g5{)ggDZl8iI7_Yi5$?@OuKb4)wDQEao;z(~LPBMxbFX4g z)bWqu!tlPPN`Hh4!&ddtp;7Nhjg?s*58u#^?3J~ma5nyziju%B9kiV-Yk7)`t%+5B zuN=5GzbxmJVfW3tFdMGCx3XVdvNs_bXDF7G#+7ex_zH%wP-@EfU>4`xVy~ON=cEW_ z{}GacD)(DC4H~CGv|Lr7%B$l^BS-|L^Zj;HyDu^wHS`r3q>6Zoqz{M%Rddsl(R?_0 zLB#X-D{#t5G!7K5oZRz)xwWh}5vChgtgoc13HTLsa*DMqjXG-`MT=Z6s3Cz;c<^(5 z&(Dk_6mUw01AfA*=#uD-vBT&0rKmi-vj_9@`ln7ql_E1wv)*ka1cUvx3(PE$0PA)r zAA+VYYjseeu{U5m5edK+okmW*t^OyFG`d>ff*Q`A5lWtzwTJfd;USP8sZkJaf$+UmO88tcMTFAMKJ8b3c4#jJjrDO zY$J3o4j#CQU#Vn)>P}`Ez!9OSN7eo4(I%)6#{c25_rbT9Du*xYp0zFY1ithZcxN{c za$P>XULW95rV~k9bS_qj#k$K_4U3}uBC~vrhNJe-sY9l;;V5T8Zf@g}GCw%R1gQd% zeDTpqPIhWa-Fo+f?oQ4--bTbqh^biELgCD%mO1=I60bjZ<3w4^Ay`Xzn(Ke){jB`F zi1!qR6dZjKyg*2A#dk$r65Z{=gQC|r!zO-}rja5XM*Ig!gHeoqlg&X46_|2Xt++J!_*QrmF~Q&RbW`HC60lFZ*tJ8nK-T$mLn`d1XDU zdyVjDUOUSndloP#F8)c#**Gy0GXBB}EVEiEsJA=I7IG}D7Yd*2q3qW+*>_kgjzUD` zJ+d4!kG0)Z=>&B(;y0-L3(^?%pWXlYBp+4%%A@{ym0_U@w00!(Gu*o;Tj^P?VT1~C zd+zzl6In4E9aeB`xPGR3;vM=tN!orY!&jNqG!`SP^QU48R6KB&gT{M*coL2ucp7DV&KC;A+2)(z&bj1-RJuF6zb z5n;nI{o{QN&@Xy|!Ux}J7%FcE;vc2Mhg8Jc`)S&WvoFNuWHcCPpJYp{?9bWOOeIL4MHuF`H9 zce)TIT0r0sAOK%z_^ER9PdHQ=Q7TDJ9%b*&I@xEYjUOFK2VH%f;Yd+D<5$dLe=c5s>I~9OI#_F5@kzXBv_0V z^Bgh^R#G1yi+j*U9G_8|nBt{-V%PLUyFkNIOx{u539b)v^dQV}3j7v)t>J=0IgE3!^G1-4jY5y<}KWE-Xt8^2|LPf2$rROAu~Mmqr-GMiqTR0upf5ek9--VhP#$wkTRTe8l;Bg1mqd5)>!D5U+J{ z$Ax+PUQM2?nOB;AHKew&>LGgbe&b{Qkx&=A>shPEn^4SCio ztMfw&$zzrCBO%46t~I@bjauch%^0-`Vp#ueb{^O)bL@LJrw1mrFHKygs~iQ9Kw1Rn zU0ic+4YPWp*2bw|)X??e3MwS)YKJp+c07{q;VKAUPVLgT#XVEl z;doo)UWXCl4z$8)^)hH}b1d#RUFE?4r{?Q_-n9LpC51XIwLEbzpE^lS#p@6_bI#sx z@z`@_XSlUcmDQ+WsP$_VR%ZG_*-f>eD!Fp<0cd8A8?!&Gjv9j@LrIo2u0~P+$sgnX z*)@0kgSTyrD<#zoZqagP{;=V_v>U5U0m_ED4lWj9-xQbEj)%|^_Q2X0t$_Z4kB#pR zmnFUX=)FP)V+S%Y-7>XgEy~S|@DpzgnIiW43A}R)e)k=%Al&e9k;{VyG{}zJ8koAr z87(l|?Koe5N*G`?b|2hFndE;~tMSjC{rhzse@haRg%=nZesRds{Z{TUY4osbMIjC@ zsQKKaikZy^Q}SQ}^KqH84ZqSES0Q(y_VoH!6Rt;lX8|(m`S~?e=(o4|qy>RKG_@qv z(s)DnXUp(qB@3t!N1<9VCF&!7K8HwO6`@a=2J0RY6y(8kQ*tX;!s}+S*fy17T@n$u zXfm`{Te4~4?ZKIJ9DVcdZw~qd`Xy3?FQu+ z&~H!C&_x1`l}VT0p<8|Oh_}in2*fyS?|dTF=QgBDh~b<}DD>jO0!%QiA|^**tD11g z&Q=x@U?`|%NBbD5Jw~HfnIIvQ7syiei8b(Rd5=OLm0U;CH?G%3d)_)fkFfE)_uh?A zj~z5|r>~L@kCO3Je)`E@I@<;Stv~-{u194-p~33dfCYkjbgm&{`hHbZW;qHA$K4Vp zr-9i+q2XaexFF9anY?-jpn6T@co|H{-pvFfU6>8z^on63HuMT7w&Uum9JCcj#FY%|o zkLBRU#pg5}3v~P}O4=_2R<0%!NMRq<>8jIsu{~A@drErcvqk6UxgNQNV4@Rju-nsT zKtK%5$mqXvBIshpN#NS3+J#5_K{|n(uj^k*-sXK3~!PQzIRo! zhl_aY-nc6~c4;*rg8#))geXklE2Jx3kJ-gcUfWFc34y}-XnI7M&ROHZjc(;BanG1D zTFm2)R^D^6WiTN$VYHWaVKPr~g>wL#wJACsOCFR>o7jUb(+=;P%3p>`A-%NGpa-mQURX>A7K;v zpdnUM^4cR-WHcMxR%MQM*5kI9gHs{_y?z#YScZqt%_Z^TmZjN);Rf#pfHafvVfb7s^FQfWq)xAtKpr%1}j4Bg|stKYt!%0>SB4 zN`ZZgL=e}&{~F9N!En8CRR@FtG#glZ!?gfz`k6ow6PA-;&5%X+lb<#qyPyK{z8HEB zFOeouCzuQ>*HF;|1{D1shuTlTrJ+KLoE)OMuE5Nb(%EUS55xIsEPICW8Stm{Kfd(J z1k=%hX&3t7ZsIz@W`CRW4bA9pu1%lH9e~JV?qsN+SQc_gv`L z_`e;&K=~>#?CiahHM1LfMi09x*h53S4NSTuDE~2E%qv_G=9w*c)T*G`uS(#}7?Ma^ z4bhKNL2#tjj7Tx{TL`MHYP4<#<@QJxYK8~BUe-y|co^=FZZ;y0_m(a4^!FhGe<&gb zLLTKRZx0LB+j#rWXy$HNiv8;$$|O(hc=?Hw!}hB z$GbuCuA;Qhg(7Du-dIpn?smB3;9t%P4e>qeiHy+C^SBt(c^E*MjK}29jMw&!=~Xlv zOOoTu0kqY>y~=f9LuD4SqCCpWZOt0Hpp$rwJ%}eld8qK=tIuXA4ogu@s(0oaZsy*M z%DS06?yCze*my32N=8f@`JAAJ%-qySp=T;#8idv?};hY+|VbYWZb{KoAcbN5;RQmHLq3@S=ig^W4;*x z-s&DcFBdSUYb5F6@GH+-Ryp@bkJDNZrVfu%!am=?vv*^;)lqD+ZdGqa045yMpZJP0 zp z6E3mJHi!FLWsK|BW7}V&zbrBOn3(_C)>VY-kvwf;sgjaK27iwoNvDZ<6+fyAt(nuR zJ-=jbRLna;N#R9N zcqy$z9c`GTK&Zi(mRdi$h^hnyk-Q*3ei3Z|m9EEiH$jx<)9qsmR^bEaOdzKs_mXx1 zrfH%E$I+A9Z|Ukv7^bLpm>pQmXZihrIX|rPkoE|3cdyh}3=Jd)auTiPI*LlVLQdX# z2oBt{&n;`CO))gQmu<}-?<+7#FO$2!Ms0nv=Eaj_!8-q>pyh6+Y{x@kfp%p8d>fV> z$ZxzJEVw$9U=V4VAGFqWfb~R6I7FZORn=#QuU6`LMi$4=gvReXn_JtnEFw`7cV#=ULE&z*)O(WQqt6P6e9+{FcnQqv{=|i9}jfk z-44#7wItK@07OYqSbvac4zZP=OmeN9^w!|n=dWt~UiQhhD<>PKbKG9MqeZFKgr~T1 z;Nl!5;`ahKu(B@c-mqECv*}^UbnB%CC8GF?fou4mKY^7(lne?L_mybAv(=5f@9DdN zM@?p%&p#vbgp_^JvmwFk-dqCe^q%e9z~Q%Tt*ZPk3rhRx6rZnN8`{4el-m+YSJ9R5 zQcIUD*HU|%EmB*7ZcU#$Li{gmiPuDEO_)RBHbpOHa8%M&EF9>{cPFR5a-}{G;BE`h z&tspHxnt7Q#E2(Che%0Zt1V@+gnDJn-{|J<BYL{L5)$|YBL1=OIuifrnCWaVpzwtlZU zU?ayc!AAe$zO<Y>JP4ml$TC--Q-sY7}pV(-<#7+NDMtu5tX3sjgueQE^nkJF*dFqO>Qff zyU+S@oXLiT7EhEpQT%W@4wkiHqM@mzzg$YLCVzc+o9Xq*YnM6B2nO>rhceNXpaD!Z zyS#6P%66m_ip-M}SDAG6{!E=TuU6%I8$agQ@D>ym1hz+9wU)UglK!pa(+4xfur z-QBwCOH9n=GWNAH4rWxv>#x-ke|JMpbYo-}WzfhiPJY;q8oqsCnoZjEekz*LgNhP0 zoABbkR)Ud6pu!N3GnHno4Aqms(S9@f#aY4v2WaTKBbE`NflBJAMQD1WwbAoN1!I~t z&0Bn2SnbofG>q_-0Bh6GrXqB9gYSXv0k^pHo`%!;Y(KhwWsW*I%~~?lD;!(V9jvsu zAyGUb!}3uhOcw*s#WyQ*xDi_7I#0*Z4TsO`7Y|MaO}r<{u<#Rz(nK%*g2|{O?vOpH ze|fJavEyNkx}KoEBYu5&4pp>q)`}U8iX(xCy``l1y-^9Vak_7%Vz2de`>Y%=%HLB9 zhr=Ic?kD?b7O6VC1{1PERqtQKwN2T1kraj5@(Z*+i8A)$lR zeWOP1h(0e|=fxDXN4Ls-^Co|5hzW1W2Vy|N6fcOVvs0>~dFu3~#>i%2$MK^Ky5==k zRXMu~(l9F1P;-Um5=zLlO#>!3qbuYpzDYs?1o7RI8XjXmEU5$RU(=gYA8;6qTmC4K7Ka!aU`AFCK z?aS=maLTd|yCwB;RTi_Q2)S^{T8~?051P1F%3wN?KMDhnqPXuumM8pkLQa=Mt|ipD zIdu{$R{Noo5Z70&W+$xMYfR!fXjYD|X8fryw08TpT7?h7pAs8PoT(P62)@1vz02>; zz%Gg=kr0R>A|)}pbKQG+`s;LWRsO5=u+T52Di`ZMVjO?+hTVVUiT;P5arT)!cW}lS zIWAUmuuHnB6GR>==@QxbRb1=uSTCGi-EnAbLqk3he`~RyLUTCvisjxVGaFiRSKhB_ k9)5OXAHD#=AOVUexz}{#BfPpmm(}anUjM~+2ITku1CzP2@c;k- literal 0 HcmV?d00001 diff --git a/images/no-cpu-performance.jpg b/images/no-cpu-performance.jpg new file mode 100644 index 0000000000000000000000000000000000000000..fe18a17475af3f49b20242ee4097b2a4cfdffe4c GIT binary patch literal 16928 zcmeHu2UJvDmhLUdl9U`Jm6AbnrVu2FNX{aXlYoFA2v~p;L^6Vk5+q0l2}O~kq{L+OGz>TBY3S$}!91*t%v>yVbZjE* zT(|iI1qB&c#iT^}C3yq{`7dt*!o$NOz$c(0B&6bJqGRI!KYpR=fa|!Rv@5qUK+M3E z>mZElAk=q&0WBaF=;sUg>kD)R0}~4y2N&-uKKcUKHQ)*e1LF!N1{M}3Ci-e0^mzdD zIu;2N-yLjH4O1LuXEJ`js4uuIvW4~JnteN90W%kWysH$HRMa##Z?Uqma|j9vi-?Me z%iWb%P*hUBr=_iZXl`L?Wo=_?=laae{kezdi&w7$0)v9zgha=@i;a8#;bVMC zYFhf&jLfX;qT&)*DZH$_qM`A7Q*%peTYLY&;Lz~M=-Bw&{KDeW@{g6(wcWk_gTtfa zlhd_6lh2M91g=*`2p4nTn8 zGqx1ptN$Gb2W@4Ptvz^|-Evf5sUcvh+LydDf|_sVA{G(XYdC23YgEta18m50$z7C%pCCBpnwy4N)$k@h`9BL0tI}DQA7cCE43%t zz{LMd>3>$~e{|`eusB=&HDU4e+Ht}239m9a@_p_l14@?CJ1mmDSn+=k>=w;-s+X$c z-mP#gR}U2A^XP!Bu2;10(K(La`$m)@Dw!{}o}c2&?=AIJHGQ-jj?E<0)CuRM1AkcH z85>XYRXVO~1J7Ee>|`MPexQJFu_yo;+hyeTW=`!CqXz|;a-x6>&UiG2q=W!$=#g>}256v%j+Bv_s*ys?fQ5S)&*q|Eg2KHZ&2ol> zhkpau88;H6)BtMdm#F^7^Zjp?0?9D6X#m z-Tow>U=29`-URl5ync-&clnDbY$=6KLdI(6jt(OC2R3;KD0bCYCZ`1<>ooEJR@go= z^GAk;46Tlu^1y&{U+%NfsWcT8k6Ara|9UbWwIyxoq4}gUj1@Ex3cW`GQ7h+qA+Gr-!ll3&=LhKzPLon|Kuk4d$y~$iXE~j zAMY4-0r#KY4WlJJ@s!)oiG6caK!yT(8PKp>pIeKFP#mw)LIE;9Xtc&ILIHbHfJ2{t zlHIP_1?T{Zpy@yXtp`3iO^^eeOTS>xivn1Bv8S|#N)Cof>^FrKZTh#Sm8b_TGTk_% zPKStpo$2qhyz~UWikteIs6U;~uQxGU8*i#f&zP_?v<%gF)Y;1wH?qyiiJSO9=gR{= zYI?5_U^-Wo7%6w-7FrLZp{qBm9fofOp#Y``Qs}5MBxqzy@XI7!F` zJrr=JC*&~98UR@&hFrWZKmiB&cIed(VFUivk*i{Idc)1+9D5A9GLb#;>mg95ma$|CeaZ2ddGHNZX=!Wjk@}%Ys3|uEkLLDs zpw4@=esu(gI(wl`gS#M*oOg`pnbPK5w9ebNW%C$UI})WW35Xs{Ms=%Nu9*6ftaA!Qmf(HJiUG z4<{*lts`3kp&MYd#wa_ZQ7<3bYTgJrB7t6j%g`3w-tRC?sAcSaUOb%^0^OV@Z$V0& z4dS+TG&u`4Wr*4Bsg4gKG=cWuyCwtwPvxz`rZ!PihTe#!7HLKA&W08uGt3;a2J@=$LUhgI43q0I1aCg zS}66rgOix?mH7qE(m{z7y?u$+hS%^X^tT+Wl#^qhgM(dirS9Y=U4^#ZfS%)-pn%nf zLpXjGQ3&dFEwI7iVFzV;SKsRSTkB6ObI4K!j)1fJS%>aW*o{xk~QZMEWa zH26}_>_pns9#DQm^U-%GfF|Z#ma~KNjbP4O`|>EFy68aTlUWjCx}kCQbe@8<^Y4?~ zDNrZ^BcgP-cxvqhriNk0m=L~AdU<5w6h>!wACFY$gAB>VJZ1BRN}K`niW(BdYQm0^ zYm!m*j@=FH;K#V?tzW`@dNAT&pJ3_94nI4AiG4CCEyU8w%!;a!CJ^GR1dwWm3y{X9j9U>8eL#@$cx^P0*;~Kou zns1eb@iyb*KePic<~>B%^$}Iy?@0QyH2a*wkfd!nEBxu2bw%%Gh~ zOd=#)C0>;gaT(rVQd zCvJHuaWe+pQ&pWdvMu$yb@4WF_Q`pj_v9>rp32-5u}N?4vngR!^ZRAriu|}6@;`!I zDwqQpAc-=P1PQePnm4Yr1UZ%@ZSo$T(dS1kpE5jeBSG%dmiK%*c5N2_yaRI8+|5sL zF!$ji5f$F;Y?B=BKPBqhnZ3;*j4LHRyC5p`QjOF$%Fe1O(V>IisHy?yc3NIo;YYp_ zzR4=<9MGVjKK#+unzCNU1Mu)LH;?`grBtK)6{oMGs^22&pVI+kiToa7UMDmRo2vy2 zUc?;sgVoMFiA4{n9B%4rJ|&CA*^1_Ws)z~DBz;%C)p4(;5+1(x93D6~k3D@i+?I*3 zn!J61#kTDSKsq`cn8%i+zLQ%|kFiEjZ=snhVEgpxofV%Pj`O*bnqJ2KPnsNs)#o2E z>5@uvJ3<{7oRUs9Hxq?~omgr-Yd{%_zvCsI=4ho@&>30_xH1S&8h zx>Jx{g(x72)9L)6cFq8gw|-{$ftAkp-dD4;oP41ae?eY ztzqQ(G-G*<9{gOgXtke1fG(oAeG6eu7x>UetjXXOLcIHv^Aq^B8JV4uHzXBh_0NIq z3nj~zLEV>M7P#@%3J13yd?DD(xC7poBM9kvJ?Bfk_yd1K35WcX^W2mvOyXVe=qEgw zhY+vOk|`rD^VE&H59`#7Tic;Tump(zVbpe#@s1p_BiL!5X~^3tU4|HcG<1I@ zdlR{8f9nSaLGa6&{;R`)vVwMUEh5v%Ze*BL`RX0kv0)0L%W zU^I>A)Reu9?|vJ&j$N&=m<=-)=ZKhY@GYMQVkkhWVuU+gEz9IYVS4%8h@wWPt7pP- zh#X%bx;2@{<=WG#{koQFx5Mtq6Y24V`>=1SvS+w?TZyaQwe0Xp2}mbBwGzR1zwq0{ z8u38!oY&ilJSr8By|xMi$v3xTqLqmjmWq9`@gBO<7lE>+NnA_q?mQNvy(oSg z#~*W)^`d2I)-`DX(htvdNG~bf`&v_0?H#i`iQDKSEw6E=>-1j`qx=Ig&J};5I(Svw zFk5;;y2Bf;Kq0f+Ro`@3w$5lecPgi(wu8rys7%YmqqoDU?`J~#9Ciaa(ssie&xV^5 zx?ZcH*A*0lJd>!+v$lcmQSFLcyhdBPGWR+Mq*1M=)>G@`Bii0{(*npeWU&5rkaw>S ze}jG3w0n`5%>DHvV`)*L^;_2{gepUKOCam=RWcV*@PaL5?EU!X(lhd-a#Lr954~NI zduk5uEI#4dN}bd?!8o)=QLm+cmgurx6mIvqQqi>lXfu@C%S0~^LRwX z4lahY&M&wzhQR z9#o(%;907$QNiX5bc$7`vj(g&Lx02!fBZG8Si1ge$_AaX`aACZk6ivg`<{+|jl;dU zN)BdXgh$dc$B<6(Y2upN_I;=f<(aP5NlFz8P|$X_{Cz*C%zni>P1T&z(OQ2vcP&#ZBoGrEC!L-a~-#-a0tU)v$js-mOCCcC$%_}+|;+~KIxG%0K*QF0wj#!kLW*gY{-YH(P zH-g8$wRkG0kV@4fb};WN#Gb!SL+O4sHOt(WeW=r;LAWyJZc(@>IL_#z{k`06;nB&+ zg_E*H;%&f79p^soV_e2pt;~hq=o~SlCkoiO2}J?vt~yXmV$M5pWB$2u$M3Hzp~IhS z-F>Pkh7^KXqR|$+hz57p;QajX4ENpy1q42Sp{EKmmaSDn^+m$$vZoOJ<1H7fMK66xFR%=AQ}dJ|@@L3&NC4w+30?b#-}fh! z$)1x;R1)BsT{Vzuo>g2XH1&9d$6es9NC0;=u2@s~_!XDFw$GtQf?O~BGpHX1UdjA} zjMOs1ou2f?ypFHfKXU#Rer8ULA&L4|?T`k}Hc!YG1Wh z!2*}*0{50f#TTE63Q3WH$YHSRgbJhada`Q@x1+xM=_r1qyV~~Og1KZvys}~vnIaf= zZb4=thA^EcH5PF%_u9LXO!p#sufhy8hb5;DQ2=j4lbVXB6D6zoG>6p`#}J<%UvQ26}Uj~Qb&)3?1W zzeTDzTlF}3aUZ~W`chdxyjD_H`#89xB(XOf~#xNcs4WTq&sMq)Jixql4L#=&j|wRWq30$1m60F_h>k z5FI5fbOo##-ow1F3`mX|Y&WU%VN$%C_|I=H7MJ0+)$lEo=(sXRz3O@C4?{SbM>Iw_ zO9~eoO|%!^E1slv2fd(5!+a9sK1mmPuppKT0ek{~s8>|H}$Cb~Yl3iM=ek8}S zVZ}^^Jf=5#RfoYyf8m4hCkvb&*xRcNF!krr;3g^? zEZSPeY2vqTxaiM6g4t)t7nbzzRHoT03mGUh;xYOb!JRRv3?rM+|ARC=A_L;v`@k+iI zsbPsXQqH4?-Ru+~!{he#klw8QBC-D@Y`b_G0d75yAiJX;v;@AZWYbp zjN)B}(L9((H0LL1Ay+y6jc%%METIQ2hJAeqiJWo!H(6`bj5Hs^$5lL1DO9j~H$oq1 ze?|d0WT5^_D)|2-tMezDf@OaMF@Huq#luXsS+)Slx2prm0FQ~VbUz)xKiWm79kF9; zs&V!qv4*){au<;%1#4RU?C6NsFq-_$7j(td{QJxiiP{Aio>%cJbfQbIf57wq5-V#{ zzK_&5=y>jO&h$owuNewRB-TAo#`}#cKk63(4-39ERuD}DmxU^k+;@CK3 z^c-^D5_1-#HK?h2$0Wb@WxCo^wEoJ<{%1J$_p7$K2~NrtLpR#;w&Nx?`cEfr5DOXm zo+{G@O-V+0-$>Slk%8XMN#Jh^Gz#m!c*|ZEc3U=+Rxtzf)nWD{(*9>dA?upP(!C}% zsV`|(WNU00?A|@D5$FYtz@8oc-R6gXHBsH!AEPHx7wv|8-|l6c8wA!V8M5(q9S9d5ou z-Zpm{yuzP-HqvxH7K4*mmKYli2A1|2(?c=oe&j#OHs5f7y6fRwyCBgwGHCJ%_ z<)as`t@hL~-_`{ys56wZXxA*p@!qrmH z`IM5M^$$+cUy{)1%I`1ciZ1>B7T5iICdR&^Ibg~^b4`Ets+*a|`*UuGJd452@okIQ ziE*|`>vvfo_ZzhSYLWzwqRlecy-e@;BkxdSKF>Az#hswbF{?4QhL9NXCfYW^7|&~; zTL5=178BFYFRVqI)OS9=cpvM=VmwZZkvX&JYjpb^159ODO#latjR3L)H1J;&CksuP zD?ERkoeGZe#UdQV(GlH!?)d|wg3Tps#i$&C!;rv4ViNvb(a!6ct>S^n;Q8~iY;xI5 zkW3v$axOG@;2e^KAYLCvGHBSOgpB%LGOeF%>({0Z`-gdrMoOhcO=OGS5nNXCI8$=$ zkPo@M`04$RPQsk50`2{>cutvuzXx4V%us1|DpGAh(}8E&e@BzZ)ztf+O&MGl;+F2; zy|zTLOdiR8NbA>x?lXu@A|95Uw8Z!br1f$I`TU^&T%Z1{553#GcTk509zFx(_x2&2~U-FK%Yd6 z-A`rRS-)@W^uS6-gjuTPYJU!M;Q^K}&Q8{k2$zd@P0wchBgkaS z-u7{#-Pjj_9J65cEDETen=@3}y3FicgM1|altWl>7Td-=M6S0<$+|Trimvp@GqlD_ zb}aNygAPBS+iL`?bG;ztXv*&bSxrG-=W~qMEkze-x=_IPzM6BS#{_g{8Qt+}bCgqj z{=tPGeThg{^MuXCu>kS|Lw7{+0~D~%3qDi$qId|&yU|yTZja^d6!R(?(z9Q;G)qz6 z63=~HjtI1c6mLe%gra+|jP+~xyY#iS)ib|nyJK95n3nE9(+9@mdowm=F&8?ipml)& zhOcU?iZO0X7n`l}$PM;ar)Xop2a~_ijRHuWb`ZIAQ&9vI-zK8XWvFZRW8xs1oBpOhI6aCHPA4~OzBpD7L=$6{uI|H7f2w0DMjF7%FS94p-~mrnv8%o7&5}Ht zSUt8;0uU=Tm1tW%YkcB_Hil>8}9Iq^<2k*N1e^1M{rpnJt0 ziKdeK_xF3)RQ;wVSn*?gMFEQ)9M(J&UES_}A|fK<7jK$6Y>l=qw0KA~rcAakXl{3B z6ujCydC4!e?TZ4&Y#}8G8R#rgl0rhrF|Cj=1dD5_s>GL{^$g*s$8i#ZDHy zbJ|GMsh;B_qsy?vqANW$uXcu2j{ZUCqTX<`v?laee9Lk3G zd(_M~EVG&Wi>1AZrN72AjJA(`BIz>7y>;9FNIJ0{shxiC(i`(>T@I~bW&QAh_(-x4 zv9AF~fP=(B4kMDw{S;Rz?6Wa%V_v(ncwI$p%(KwTz|*G#5iTF|h~v5KdlxTyBP%Gm zj2&;$nrdevNlD~ckJpc+t^oIET1QHc`k|fOR@AT)$|p2Z=@p5s*@`RkUs^(L)|Pqt zbLU1j`wNxoHC{OnCm5wBsM`Nbge~n@^kLv`;TxK}xXyuGw_OL~bCjQ;04$_}yXq9u z3DNz4kYaDzoJ-4Ycrnz|y~Byl#P}nNVBFcuT=dGs$1Z~R);%7Zq5zO$J)-U?zab<5 zFT+qwB_)qMy%{>yhsP9vMGdz)eN#B=d8ZRRN8 ziZ1#V;UlW@dfVb?bGnJV>oTxP>pj*zwI-g|<7qYrj@&~h@|L{^P9=u0J#_CUkMvT4H04I)8$8FW8`{M~7m zL2UOECk?0bA4$U%GzyQi4T7&?E3^8td!eEB^5zA$#ET{or=J6A%JcHN7CKVaQpCOz zlLVVX*=3nGFeT)Jc zlpz=RFJ>;z(9AwuN21G1D{fuQ%K^tAoGl-JGAo>KS#e3}i(=2kk=(=FqL6KroT})A!{`_lnp#eRtJOB3BpnvwU z5hu#5=z71C5!#(-KzG}x)4+f-BWN&Q%;f2E#NOwU_@l4mT_vl@>&$K3%gSbZ3mr;b z)C$eWdJB5!+dZ^X`bAtLGn-Ipn^?(Q>z8v8-Jk!OG7i_96iY>XrTlFz7vKwLD!xnXBUIeTLs=m62mVp3oiUgPoMTetWvjiP@;Vulb6Sc7T~&Vi3-k z{Oov(P$d?gFQXRhl1UAJeH_Do_P8=|XsbEVR~ddiNxy>(jlb+V8V-L^tN(4+GyiSP zf8Xi{JM)v0XI+a&uH*QXb6n2dfKLjp7At0W=G40Fl)3Cyy|x?IQ!LsMtBXb}fl5j9 z;dr%j9_74iYiu|yEa~KT1O84(OZ|wcrt(|A!N%iwqc?PSzp{k}avi#)Mw3}Al-HrN zS4lt>i&8Z_Py(jp^aGS1^-B2R=jbV;k?D#G3-uykIv*AZ@fGz~nC5tNO61yJpt>xZOM;Fe zVO&IcCe=wSxyz)$y;##Qi?J6oh%rMYeb}~gR0?08Q!&0rl>xiKu+I|WMi8)zYyPU9 z>4DTM1})P`l3?2P7EWf51rN^K;jQ@A8tSR4N!K5tX%{hb~%GjoVun_qkQFQFV|dG6zz} zuPHok@4)tM9Lxe=5p4xOjZN*|{<= zDV8G*IX{oMX<;J2&6TA~Q?)#KZJVL+Zf)_?O$+M>Y9HjTWRIFwb7mO$H7>dy!W2b zIZ`|-rgjcs8QA_3%mxe6C<`Cuozibi#QEk>cUe%YrVJ7k<+gwYDGV&YV_zwg7^b8~ z^+sajk-watrd4?cFAWqm4RugjP|r^71=S_vU6O3h(4mqJdM%^)o(C|+LBdXfX*(W% zhbx;*8WBwGnGNB9nVQL`;@qq%f#iqAu*tI*j&+ni2es%JNkt;b@}^==?D_BU${Ob^caa`91?LgZQ}r1;aNXB zY%wmK;g){V`-zmiIS-pVg`bJ5xsTMAan`h+Zg(a_aN+`DL=&H4I3pBU6e1d!LvHfQHCBgfZb5tS= zW)A5px-~m<`@!2s{a^k2Evkc?p5yo6fA;Om#QMaV421rmw|`T)lUOVjKPm56XlzFb zXZOU7btO!%r`;QCygizMAsG{T9OX@!rRFd%L3zC%!${ROt#nrLJ!eVYXBX^ytbOSI zCoQM>E|O)=H|Et1cQ787M^i@d@#-13i)0=?45|FOHf-2b*<;EV2X>|Q2~Ca!EihU< z4p+Eym`}0YIG&hc=?*D!W80!ssqh!@$xQR@0DH+QkO$Wz_Z?P554Krg=BasKq_t?B zDH$5v7Y>rI>b_*}E7*?fA3u=H*H`d{YriW{4_6pzCZ1wV*WBp|&e$bN&5qrRYwvf( zR9;jnqJxqOJa+AA$5y1yt>hNa|D?7v9VQ(;oRg-4-{oN$mg-DGenX$Yh`VK4r;fG2 zXPZd>;V7ZeCZFfam+6ALh`KF>W@4Q11Y$vuocai?T8vhstPmFQz^JKIpL$}GD2y>w z3HQ`Rd6JS|Vi|U@`9LL1OLHKfoMOIjXz;6y$3v2^9`8V{t?~uYd{ZcV3!yYJUJ&Ty z$OQlV7B_l%raj<=XBsts|Mi&G(>rz4iyQjv<-5xT9 zegAB`(LeDy&?9?ViMN?N88CmNx8A4xRgX(N*Gm!=oX2zjCJU // std::sort #include // std::vector #include "testing_helpers.hpp" +#include int main(int argc, char* argv[]) { const int SIZE = 1 << 8; @@ -124,7 +125,7 @@ int main(int argc, char* argv[]) { printf("\n"); printf("*****************************\n"); - printf("** RADIX SORT TEST **\n"); + printf("** RADIX SORT (Single block/tile) TEST **\n"); printf("*****************************\n"); zeroArray(SIZE, c); @@ -143,4 +144,37 @@ int main(int argc, char* argv[]) { printf("Std sort:\n"); printArray(count, d, true); printCmpLenResult(count, count, d, c); + + printf("\n"); + printf("*****************************\n"); + printf("** SCAN PERFORMANCE **\n"); + printf("*****************************\n"); + + for (int s = 4; s < 25; s++){ + int ssize = 1 << s; + int *u, *v; + u = (int *)malloc(ssize*sizeof(int)); + v = (int *)malloc(ssize*sizeof(int)); + printf("==== Array size: %d ====\n", ssize); + genArray(ssize, u, 50); + + zeroArray(ssize, v); + auto start = std::chrono::high_resolution_clock::now(); + StreamCompaction::CPU::scan(ssize, v, u); + auto end = std::chrono::high_resolution_clock::now(); + double diff = std::chrono::duration_cast(end-start).count(); + printf("CPU scan: %f\n", diff); + + zeroArray(ssize, v); + StreamCompaction::Naive::scan(ssize, v, u); + + zeroArray(ssize, v); + StreamCompaction::Efficient::scan(ssize, v, u); + + zeroArray(ssize, v); + StreamCompaction::Thrust::scan(ssize, v, u); + free(u); + free(v); + printf("\n"); + } } diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu index 3f4fb58..27bc8b8 100644 --- a/stream_compaction/cpu.cu +++ b/stream_compaction/cpu.cu @@ -52,7 +52,9 @@ int compactWithScan(int n, int *odata, const int *idata) { } int *d, *e; d = (int *)malloc(n * sizeof(int)); + scan(n, d, c); + for (int i = 0; i < n; i++){ if (c[i] == 1){ odata[d[i]] = idata[i]; diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 539d98f..6ca121a 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -58,16 +58,35 @@ void scan(int n, int *odata, const int *idata) { int *dev_pidata; cudaMalloc((void **)&dev_pidata, m*sizeof(int)); cudaMemcpy(dev_pidata, pidata, m*sizeof(int), cudaMemcpyHostToDevice); + + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + // Scan + cudaEventRecord(start); for (int d = 0; d < ilog2ceil(m); d++){ - scanUp << <1, m>> >(d, dev_pidata); + scanUp << <2, (int)m/2>> >(d, dev_pidata); } + cudaEventRecord(stop); + cudaEventSynchronize(stop); + float ms1 = 0; + cudaEventElapsedTime(&ms1, start, stop); + cudaMemcpy(pidata, dev_pidata, m*sizeof(int), cudaMemcpyDeviceToHost); pidata[m - 1] = 0; cudaMemcpy(dev_pidata, pidata, m*sizeof(int), cudaMemcpyHostToDevice); + + cudaEventRecord(start); for (int d = ilog2ceil(m)-1; d >=0; d--){ - scanDown<<<1, m>>>(d, dev_pidata); + scanDown<<<2, (int)m/2>>>(d, dev_pidata); } + cudaEventRecord(stop); + cudaEventSynchronize(stop); + float ms2 = 0; + cudaEventElapsedTime(&ms2, start, stop); + printf("Work-efficient scan: %f\n", (ms1+ms2)); + cudaMemcpy(pidata, dev_pidata, m*sizeof(int), cudaMemcpyDeviceToHost); for (int i = 0; i < n; i++){ odata[i] = pidata[i]; diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 853b719..310a799 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -33,10 +33,23 @@ void scan(int n, int *odata, const int *idata) { int *dev_pidata; cudaMalloc((void **)&dev_pidata, m*sizeof(int)); cudaMemcpy(dev_pidata, pidata, m*sizeof(int), cudaMemcpyHostToDevice); + + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + // Scan + cudaEventRecord(start); for (int d = 1; d <= ilog2ceil(m); d++){ - scanCol<<<1, m>>>(d, dev_pidata); + scanCol<<<4, m/4>>>(d, dev_pidata); } + cudaEventRecord(stop); + cudaEventSynchronize(stop); + + float msAdd = 0; + cudaEventElapsedTime(&msAdd, start, stop); + printf("Naive scan: %f\n", msAdd); + cudaMemcpy(pidata, dev_pidata, m*sizeof(int), cudaMemcpyDeviceToHost); odata[0] = 0; for (int i = 1; i < n; i++){ diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu index 1775000..5eb5d66 100644 --- a/stream_compaction/thrust.cu +++ b/stream_compaction/thrust.cu @@ -24,7 +24,20 @@ void scan(int n, int *odata, const int *idata) { hs_in[i] = idata[i]; } dv_in = hs_in; + + cudaEvent_t start, stop; + cudaEventCreate(&start); + cudaEventCreate(&stop); + + cudaEventRecord(start); thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin()); + cudaEventRecord(stop); + cudaEventSynchronize(stop); + + float msAdd = 0; + cudaEventElapsedTime(&msAdd, start, stop); + printf("Thrust scan: %f\n", msAdd); + hs_out = dv_out; for (int i = 0; i < n; i++){ odata[i] = hs_out[i]; From 5aee3b5ca826620c4b02aa3acd9b2cb680f20aae Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Sat, 12 Sep 2015 02:41:43 -0400 Subject: [PATCH 6/8] Update README.md --- README.md | 44 +++++++++++++++++++++++--------------------- 1 file changed, 23 insertions(+), 21 deletions(-) diff --git a/README.md b/README.md index 55c615e..93227b3 100644 --- a/README.md +++ b/README.md @@ -6,41 +6,43 @@ CUDA Stream Compaction * Tongbo Sui * Tested on: Windows 10, i5-3320M @ 2.60GHz 8GB, NVS 5400M 2GB (Personal) -## Project feature +## Project description -* Add a description of this project including a list of its features +* Scan: + * CPU scan & compaction + * Naive scan + * Work-efficient scan & compaction -* Radix sort +* Radix sort: + * A single block/tile implementation of radix sort. Splitting and merging are not implemented + * Call the function as: `count = RadixSort::sort(SIZE, c, a, 8);` + * The function will simply write sorted result to the result array, and return the array size. -``` -count = RadixSort::sort(SIZE, c, a, 8); -``` - -### Questions +### Optimize block size +Block size is optimized based on fixed array size (256) and number of blocks. `[1,2,4,8,16]` blocks are tested on each implementation + * Naive scan: 4 blocks / 64 threads has the best execution time + * Work efficient scan: 2 blocks / 128 threads has the best execution time -* Roughly optimize the block sizes of each of your implementations for minimal - run time on your GPU. - * Naive scan: - * 1:m - 1:256 - * 2:m/2 - 2:128 - * 4:m/4 - 4:64 !!! - * 8:m/8 - 8:32 - * 16:m/16 - 16:16 - * Work efficient scan: 2:m/2 0.1117824 !!! +### Performance -* Performance * All performances + ![](images/all-performance.jpg) * Exclude CPU + ![](images/no-cpu-performance.jpg) * Naive vs. Work-efficient + ![](images/custom-only-performance.jpg) -* Write a brief explanation of the phenomena you see here. - * Can you find the performance bottlenecks? Is it memory I/O? Computation? Is - it different for each implementation? +* Observations + * CPU profiling is very unstable, as C++11 `chrono` is not able to return reliable numbers from time to time. However, even under this situation it is obvious that CPU scan execution time increases dramatically as the array size grows. We would assume CPU implementation is `O(n)` and therefore array size is the performance bottleneck. + * Theoretically work-efficient scan should be faster than a naive scan. However somehow this is not the case in performance testing with small array sizes. A profile report shows that up-sweep process is generating 2 times memory transactions as a naive scan. Adding down-sweep process, work-efficient method has about 2.5x as naive method. Also in general naive implementation has a better Occupancy Per SM rate, in term of number of its active blocks against device limit. All these could contribute to a faster execution of naive scan, compared to current work-efficient implementation. + * For the sudden performance "improvement" around `2^10`, the reason is simply because all of the available hardware resources are saturated. Since the block size is calculated based on array size over a fixed grid size, the current naive and work-efficient implementation does not handle sizes that go beyond the device limit. Namely, for naive it is `4*512=2^11` and for work-efficient `2*512=2^10`. If they were to have some routines to handle bigger sizes, the performance is expected to be similar to thrust, because the array has to be broken into tiles for calculation, and then merged back again. + * Thrust seems to have the best performance. As array size grows its performance goes down, and seems to be worse than naive and work-efficient methods. However as stated above, if these two methods can handle bigger sizes, thrust would still have the best performance among all. From Nsight timeline we can see that it is splitting the problem into tiles and doing work-efficient scan. Thrust runs in 40 blocks with 128 threads each, which accounts for a size of 5120, which is about `2^12`. When array size grows big, building tiles and merging results would be the major bottleneck. + * All of the GPU implementations are much better than a CPU one, in terms of scaling to big data set. It seems that for bigger problems, computation power is more of a serious problem than memory I/O, when such I/O are constant. * Test program output * The extra output `XXX scan: 0.0000000` is used for profiling the execution time of each implementation From 192c8a58f0b2411250ea9bb6ce1c9382a9b9fc05 Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Sat, 12 Sep 2015 03:31:58 -0400 Subject: [PATCH 7/8] update write up --- README.md | 163 ++++++++++++----------------- images/all-performance.jpg | Bin 18519 -> 17286 bytes images/custom-only-performance.jpg | Bin 19431 -> 16755 bytes images/no-cpu-performance.jpg | Bin 16928 -> 17270 bytes src/main.cpp | 4 +- stream_compaction/efficient.cu | 4 +- stream_compaction/naive.cu | 2 +- 7 files changed, 71 insertions(+), 102 deletions(-) diff --git a/README.md b/README.md index 93227b3..384a224 100644 --- a/README.md +++ b/README.md @@ -24,6 +24,7 @@ Block size is optimized based on fixed array size (256) and number of blocks. `[ * Work efficient scan: 2 blocks / 128 threads has the best execution time ### Performance +All 4 implementations are ran on various array sizes from `2^4` to `2^18`. Due to memory constraints bigger array sizes are not available. * All performances @@ -38,16 +39,20 @@ Block size is optimized based on fixed array size (256) and number of blocks. `[ ![](images/custom-only-performance.jpg) * Observations - * CPU profiling is very unstable, as C++11 `chrono` is not able to return reliable numbers from time to time. However, even under this situation it is obvious that CPU scan execution time increases dramatically as the array size grows. We would assume CPU implementation is `O(n)` and therefore array size is the performance bottleneck. - * Theoretically work-efficient scan should be faster than a naive scan. However somehow this is not the case in performance testing with small array sizes. A profile report shows that up-sweep process is generating 2 times memory transactions as a naive scan. Adding down-sweep process, work-efficient method has about 2.5x as naive method. Also in general naive implementation has a better Occupancy Per SM rate, in term of number of its active blocks against device limit. All these could contribute to a faster execution of naive scan, compared to current work-efficient implementation. - * For the sudden performance "improvement" around `2^10`, the reason is simply because all of the available hardware resources are saturated. Since the block size is calculated based on array size over a fixed grid size, the current naive and work-efficient implementation does not handle sizes that go beyond the device limit. Namely, for naive it is `4*512=2^11` and for work-efficient `2*512=2^10`. If they were to have some routines to handle bigger sizes, the performance is expected to be similar to thrust, because the array has to be broken into tiles for calculation, and then merged back again. - * Thrust seems to have the best performance. As array size grows its performance goes down, and seems to be worse than naive and work-efficient methods. However as stated above, if these two methods can handle bigger sizes, thrust would still have the best performance among all. From Nsight timeline we can see that it is splitting the problem into tiles and doing work-efficient scan. Thrust runs in 40 blocks with 128 threads each, which accounts for a size of 5120, which is about `2^12`. When array size grows big, building tiles and merging results would be the major bottleneck. - * All of the GPU implementations are much better than a CPU one, in terms of scaling to big data set. It seems that for bigger problems, computation power is more of a serious problem than memory I/O, when such I/O are constant. + * Theoretically work-efficient scan should be faster than a naive scan. However somehow this is not the case in performance testing with small array sizes. A profile report shows the following which might contribute to the result: + * Up-sweep is generating 100% more memory transactions than a naive scan. Adding down-sweep process, work-efficient method has about 150% more than naive method. + * Up-sweep is executing 50% more integer operations than naive method. + * In general work-efficient method has 8% divergence, while naive method has none. + * In general naive implementation has a better Occupancy Per SM rate, in term of number of its active blocks against device limit. All these could contribute to a faster execution of naive scan, compared to current work-efficient implementation. + * Thrust seems to have the best performance among GPU implementations. From Nsight timeline we can see that it is splitting the problem into tiles and doing work-efficient scan. Thrust runs in 40 blocks with 128 threads each, which accounts for a size of 5120, which is about `2^12`. When array size grows big, building tiles and merging results would be a major bottleneck. + * As array size grows, GPU hardware resources get saturated. More blocks can only be scheduled sequentially, thereby reducing the performance significantly. On the other hand, issue efficiency drops, resulting in wasted cycles. + * With almost 0% cache hit rate, as array size increases, more and more time is needed for memory I/O operations. + * With all of the factors above combined, the CPU implementation is actually much faster than GPU in the final result. C++11 `chrono` library was only able to pick up the execution time in milliseconds when the CPU code is ran in another nested for-loop for about `2^8` times. The guess is that CPU can still have the best performance up to a point, where linear computation cost outweighs memory I/O time. * Test program output * The extra output `XXX scan: 0.0000000` is used for profiling the execution time of each implementation * Radix sort test is also printed, where the test is carried out on the array generated in stream compaction test, and compared against standard C++ vector sort - * Performance comparision is also printed. All 4 implementations are ran on various array sizes from `2^4` to `2^24` + * Performance comparision is also printed. All 4 implementations are ran on various array sizes from `2^4` to `2^18` ``` **************** @@ -60,7 +65,7 @@ Block size is optimized based on fixed array size (256) and number of blocks. `[ [ 0 38 57 95 132 137 184 199 234 234 246 249 249 ... 6146 6190 ] passed ==== naive scan, power-of-two ==== -Naive scan: 0.042400 +Naive scan: 0.042304 passed ==== naive scan, non-power-of-two ==== Naive scan: 0.041088 @@ -69,13 +74,13 @@ Naive scan: 0.041088 Work-efficient scan: 0.127552 passed ==== work-efficient scan, non-power-of-two ==== -Work-efficient scan: 0.126624 +Work-efficient scan: 0.126304 passed ==== thrust scan, power-of-two ==== -Thrust scan: 3.615200 +Thrust scan: 3.756288 passed ==== thrust scan, non-power-of-two ==== -Thrust scan: 0.016640 +Thrust scan: 0.016544 passed ***************************** @@ -92,23 +97,23 @@ Thrust scan: 0.016640 [ 2 3 2 1 3 1 1 1 2 1 2 1 1 ... 2 1 ] passed ==== work-efficient compact, power-of-two ==== -Work-efficient scan: 0.125120 +Work-efficient scan: 0.125088 passed ==== work-efficient compact, non-power-of-two ==== -Work-efficient scan: 0.125248 +Work-efficient scan: 0.125120 passed ***************************** ** RADIX SORT (Single block/tile) TEST ** ***************************** -Thrust scan: 0.016576 +Thrust scan: 0.016512 +Thrust scan: 0.015072 Thrust scan: 0.014912 -Thrust scan: 0.014976 -Thrust scan: 0.015008 Thrust scan: 0.015008 Thrust scan: 0.014912 -Thrust scan: 0.014912 +Thrust scan: 0.014944 Thrust scan: 0.015008 +Thrust scan: 0.019840 Radix sort: [ 0 0 0 0 0 0 0 0 0 0 0 0 0 ... 3 3 ] Std sort: @@ -120,129 +125,93 @@ Std sort: ***************************** ==== Array size: 16 ==== CPU scan: 0.000000 -Naive scan: 0.020928 -Work-efficient scan: 0.060864 -Thrust scan: 0.017600 +Naive scan: 0.001536 +Work-efficient scan: 0.003136 +Thrust scan: 0.016672 ==== Array size: 32 ==== CPU scan: 0.000000 -Naive scan: 0.024800 -Work-efficient scan: 0.074144 -Thrust scan: 0.016416 +Naive scan: 0.001536 +Work-efficient scan: 0.003040 +Thrust scan: 0.016512 ==== Array size: 64 ==== CPU scan: 0.000000 -Naive scan: 0.029664 -Work-efficient scan: 0.087968 -Thrust scan: 0.016736 +Naive scan: 0.029120 +Work-efficient scan: 0.003104 +Thrust scan: 0.016416 ==== Array size: 128 ==== CPU scan: 0.000000 -Naive scan: 0.033600 -Work-efficient scan: 0.101184 -Thrust scan: 0.016480 +Naive scan: 0.033376 +Work-efficient scan: 0.111296 +Thrust scan: 0.016512 ==== Array size: 256 ==== CPU scan: 0.000000 -Naive scan: 0.041376 -Work-efficient scan: 0.126336 -Thrust scan: 0.016608 +Naive scan: 0.041056 +Work-efficient scan: 0.126304 +Thrust scan: 0.016512 ==== Array size: 512 ==== CPU scan: 0.000000 -Naive scan: 0.062336 -Work-efficient scan: 0.193728 -Thrust scan: 0.016672 +Naive scan: 0.063456 +Work-efficient scan: 0.194720 +Thrust scan: 0.016736 ==== Array size: 1024 ==== CPU scan: 0.000000 -Naive scan: 0.110176 -Work-efficient scan: 0.346016 -Thrust scan: 0.017088 +Naive scan: 0.112128 +Work-efficient scan: 0.352768 +Thrust scan: 0.017056 ==== Array size: 2048 ==== CPU scan: 0.000000 -Naive scan: 0.219840 -Work-efficient scan: 0.003296 -Thrust scan: 0.020192 +Naive scan: 0.216704 +Work-efficient scan: 0.685760 +Thrust scan: 0.020064 ==== Array size: 4096 ==== CPU scan: 0.000000 -Naive scan: 0.001504 -Work-efficient scan: 0.003008 -Thrust scan: 0.025056 +Naive scan: 0.439424 +Work-efficient scan: 1.365280 +Thrust scan: 0.024992 ==== Array size: 8192 ==== CPU scan: 0.000000 -Naive scan: 0.001504 -Work-efficient scan: 0.003040 -Thrust scan: 0.038944 +Naive scan: 0.915264 +Work-efficient scan: 2.785376 +Thrust scan: 0.038912 ==== Array size: 16384 ==== CPU scan: 0.000000 -Naive scan: 0.001536 -Work-efficient scan: 0.003040 -Thrust scan: 0.067168 +Naive scan: 1.931072 +Work-efficient scan: 5.762080 +Thrust scan: 0.067264 ==== Array size: 32768 ==== CPU scan: 0.000000 -Naive scan: 0.001600 -Work-efficient scan: 0.003072 -Thrust scan: 0.363552 +Naive scan: 4.107552 +Work-efficient scan: 11.998688 +Thrust scan: 0.393440 ==== Array size: 65536 ==== CPU scan: 0.000000 -Naive scan: 0.001888 -Work-efficient scan: 0.003136 -Thrust scan: 0.449568 +Naive scan: 8.706624 +Work-efficient scan: 24.970079 +Thrust scan: 0.439232 ==== Array size: 131072 ==== CPU scan: 0.000000 -Naive scan: 0.001536 -Work-efficient scan: 0.003008 -Thrust scan: 0.446400 +Naive scan: 18.445057 +Work-efficient scan: 51.995682 +Thrust scan: 0.477088 ==== Array size: 262144 ==== CPU scan: 0.000000 -Naive scan: 0.001536 -Work-efficient scan: 0.003040 -Thrust scan: 0.649120 - -==== Array size: 524288 ==== -CPU scan: 0.000000 -Naive scan: 0.001696 -Work-efficient scan: 0.003008 -Thrust scan: 0.912192 - -==== Array size: 1048576 ==== -CPU scan: 15630.000000 -Naive scan: 0.001536 -Work-efficient scan: 0.003072 -Thrust scan: 1.351296 - -==== Array size: 2097152 ==== -CPU scan: 15627.000000 -Naive scan: 0.001536 -Work-efficient scan: 0.003168 -Thrust scan: 2.192576 - -==== Array size: 4194304 ==== -CPU scan: 0.000000 -Naive scan: 0.001792 -Work-efficient scan: 0.003072 -Thrust scan: 4.034496 - -==== Array size: 8388608 ==== -CPU scan: 31247.000000 -Naive scan: 0.001536 -Work-efficient scan: 0.003040 -Thrust scan: 7.794240 - -==== Array size: 16777216 ==== -CPU scan: 46879.000000 -Naive scan: 0.001504 -Work-efficient scan: 0.003104 -Thrust scan: 15.012032 +Naive scan: 38.985985 +Work-efficient scan: 108.208801 +Thrust scan: 2.610464 ``` ## Submit diff --git a/images/all-performance.jpg b/images/all-performance.jpg index 56af743dd6bca505012cc84a16e17f901938509f..f37543b5265c263ff8820fe2bc0d5c57c5d63599 100644 GIT binary patch delta 10969 zcmb_?XIK>7wq}u|AP54IgOW2SNN7;Wl5-YXBxi|&MNkkVH$igFIcHEL2q+*RInzka z8Jg~Czccrq^PRc($K078tE;P??p=G=e%D&>dRI-xV1$HXRCQs(UXbDcKNt>J@NI=7 zL!S6au_dL-ooE?x_L%^BIrl_q){gNqQJV5PADkcU8}(5OwwQ5*yMqYxsf6+FdED^PF<(#vAQ?*Wsw2e^+gYMpSK zFK_f|COKLPF&zo?MsiHZf?797z1> zw7=hWk)cHM%r!pJI6X4-l9rC_o8(CmZym4N!Yye0Uw zxO_4F^<|Xrq_G@ewei?t#v^UfPB67JHH_~S8M&t6tb{BRIq-Q1*ENkK0S_i1$_3&` z8LO)r%Xf^u^o-=&g$Jg~D~D<|Gp3v5gaaJ^Vue11`D+uI5)zy>O~OT3`*$c}myqOl zE_4StvM|H{X@stun+&z?#dwy-vtKcuyI?!AVqIJo9+k#x%HUA^(a^P=Ae~5{TzY#Cea#B&dTNCNBt;>JyCx{$ zu%#Ba(Am|QXN!*&jABwr-p%{4L=4B%P+!*PtoX@ zLEMeZ%2#;hiEe4wPJq*st#rm=_!-oh37*iSI1KC}EV`vQ@ox-b9Z&Cn9n6(fDN(ZI z(bF@|!RZgUE#x*uTOAlXr`H}Da%#`v@OrtDbSITvol9&N3^$Rqjf`z;{l@(n51$Al zczpe-+`&ynG({(yu!)kFyZibjvw=^oWU71rlJwkqn&+mL;S=>cChWndyf(%?zh_KL zpH#zlEhxWs6I{K++p;gHPzLeMl_(cdQrc8LPkZ~p&xbv1a;1TuVqnmdy z6rjYr?NluSxg_>ZMRqBnfV{*Lj7C^2D;Ri3%zbRTW(v^OWxWRtmw0)RCUkbZjGjDGCYhByDIm;0^S*>5ffb)UG@ zi)IGCjthwTUiY*}<(TF9zE5lPtZuEtUDn?7UotolTpVk6Ke!;jni=aVSlC*Sh?M3Q z6BWT}YfT6s7D}^ri;-a?F1_tb{L0=8OjgxEb_O9mU)$fFoK^~*G*^X|9Zrc=rA82~LcVpyz zT@Caz$TRl4E=#gqPA+X$6RL1*H(>YpT_1CUHns`q}pZ}Q9x2YnM!z=r38A#w66 zA8`6en~i%JY9e1~Gd~0{k$%Fe&I#uL{#co+JMt+54GO42Z|&y&9><9X%_x95nivI4 zp*Qy=ZYAL1Hn(CNfpdN0J7DF6tx3w@DdS0x9XSQL8UUXZUlx!X{y-feyv51Puxszs zcRArU;c8ZSa-(ccWUVAeq4(}R_%8wJui*4hiJUsakX}nf>-%>XfL^IkQ-vpcLdmp$Hf!7iLk3P37;UbHWmzG1RBn+N zD{j$+Hd5vo{a}<)018n4U3lTvyD~JeWn2^?Y-T$6dL)q;O1#JwCpH8O5zq{?munnk zzzWLr?-|8yo(0{U(jH7cp*&W+*2%>=%m^T9FTyH>F$>6nD(06A9_H7$)8NKz$!yu* zPJ!8+oRo=$rC<>g*wxp3WcwhAljkl-ZRtOfdziTrLuVbblGl1*g}8e!F(LmZtjvS{A+JnIb%q-c%w7A?GpTSsKh8wX)z>MjF>XN|_X5#R_K0kQDA?$%F>&%I zjpHk&T$28F68SYTR_SiC3w_q2iwAFRBs6Mlq0>HPvx^c0rY2orhN49iacH+S^%>lA z{AZh#tFB>(Qi{dzsPhf$*aUkMdV+w)wg3HD~_=oF`*ALb*;>+6` z!YuDNL66pLkE!rzmJT?u@;VOnVD|J!J)pJduQaD487E5o+4Dg5z=((|>y##R5Md~4 zfo|O$^qa!OVp@z7!j)27r*}1pJbza|7)`X`?F-rItOY;qCDInt&nO5ba>v%fA>wnN z8(haS9>Bg#cFraVA%ELTM5^)1v?YxdtLD=>$2G~nfYKUN?Z$X61=!m$h#j0|T7fPe z5}`ct}j+UpLCs6{FnZt4bE>hip#S3GtYU zOb4=`sn5$QiKdM<&pTM_YZ^jk=0)#MQwYA3`$1FQEQS@>&wf~zTdo9=yF%te-qTZ3 zXd1=A0;xgpekCIMapg_0L+6R(FUN_$odS+Ob@pEmzB$)RQK{z=ko?$e!T3R*24Po; zU~MozUerbbU-E}h09`?uDAXvt9|gSg97Mj^0xv|+QFb*R?2J@Zd8AIos{ydNiY`ilZj@|+=ciDJ;q783Oll~5;W4GrwNA~LgEPNV1sPLqSbh`h zZ;mW-C9I6#LSHze-1zQ3``GU5Af8#HhW2o8c$lT)?I(o0bJS$jO|uasq;`GH2O;IW z9j*4xx$jj5i7tXcQ~&hnROSKvwdWGt4{6U!V*`ie0N=QpmJlr(#cxb-G8GL&=#GHp z&vLxPrbfr+570eCUM0O4o6ZTKDMtb5{d+=d*uH~)*fD(};0_xKxKeVJxyQafKJFP# zuKp{U@rs~qmlT~6AD~lW9JqO3pxtH$ajx&?xp15DMeY^MaE1{PBB_c3EYN=*Rq-82 zEFNqH*DhHK+ap^za_ZS&7@X9%x$ksk9^qvoyY$g_&=+-tVGHPjXsXTqhiYq$i4v)} zOSNezCOvd~bkz_Wi1YF}ST4RjqW6meBkar^O`}Gj3TW);bRuFErv1<`l8t0UGhK8N zJd$(#DMOuUb*O%WD*_@YRZu`bg?~~H`>Pob&D^dB7QAIGjkHbvf)y=$HJhOunjd7MNoC6fkVp;Nq{ z3$tJ>&P`1Qpsmx#!?cO4Ocz#fIlPrH`pv@lS6iaN^Do#TENs@G))yf~t3JH6@I59& ztp(8xxj2n1iIQF7DR>|(xAy*>I!susLJtH<9JsyL1 zT2_L1xbHlS2oE{ug>Bqt1dZYxf1;h~k3;>3=_!PDbV)*7W+IMdoyutmUTBT-1jz{| z0CWupP1NF2rDn~;^n>sVVU747bQpd$*04-ApO48DFRdo`Iq>cWU`cUc;e9O_H8oZX zu+3l(%gK0CkaE>y8LkA^pJ$5z)`8=q^0k1F zU*5^HTYYAvPWDw<#QaoWM2ABcsB@FYhY8|8r{*hmha0Y9rEf<>V0wSUUTH4nHvTAO z>v6w+oKIw#<=a;dX>(h#iziJ*` z*dh@Vs2jZA-U6$drAFX_XLj|v4$}6{kC48t(*9v$H-d)hhCRNWW2c;cJc+;O`F8Vy z^AV3pR{kh&UQ0`qdXs6lQF`ZZ(EYFW z2Q({oEEE11$#R7q^I*);>+Xih>5x8gVS}4IB(&X5PyfTI_-lW{l`}HZPv!Tx)fA(o z_KMpN5!|h=o1Z14>M)+8Q?U8#ZCbv_!oJ`R8+hC%3RvRO6CiD7->~mZBq&w(i#dQ0_{5vx!PdLW1Oav$0&GK-i ztqcy#iTYW9X5a0p2rJ?gS8dMZYd zXXcX|+7}xCEga21>_??%&6^M@?se?EqA_-ODS>@|h~zeIzu3ktJ3Xoh$NY|i`qTV2 zXUU;rWK`UT?}P1L%$9@o2AFy4*^&(JuvovE(f>N}H?{U>>`L_9E2~v1lwPXmwx zy-Upuzh?AlvJ5t>wkGZklLjWKq~Z3A))yZkwKJQUp&4l;h)L^zEDO|4S|8^fyAD5M ztR#e-u-=mISCPXTrZoKk|E z#4Xgb&HgfLl(|u2jXT79TOSj5gub#@Ws#UaJ7_4iu9#j>&-YgArKU+Y=dLq*1DLT>hl2s2a%<#K_pChC%;@wcWi+7XOgqp6FVBQcOyQ z^|go`Z(+<#Et6oafF5-}4g9^74c1RtVhrrtKG=Hv?;&tf$j!S-bPZZz7YZYmrDZ04 zx3f-WYD%xRoBcTQ;;xnDQPK^UW{juv`C99hs|>vqY?w#}?Rd-*I_`sT<3Ua&Ht!|z zsN;Q4k;?2_(LX5SP{6Q~X-BRYV-}@dcKwI1w!{i0S$Nw~c6Q-PMQrDE>F_QXGutz8 z1eWvf553I%Zhg~158O+y(u-o2>H6_=i(?}N9bGU20wx_IUwZ1%a>9ahPht~jdf&Wi zRR*HGnVSQbn!@gf$kwmMKIJBTLaM1N|1qC8tookSr#jxY&rIe#4esepz3vN@nlfaw zvtMnU63<5Zhy^2y>Fg+Vd;u}a*k$!gEHq8E;6?#A+$jyPDjq0ok%t7-HAVp`b7&?a z#r6zAKyxFEMluQzy~*=C>`?#ahytd!xfZcZPp(h67|HCyIm(p?yZ_3y|Cwt4O*-1Z za{3n)6Z6ma=0z%Ry}|AV&UHftcP}j50I9*xojyu9@eKH&`P1E~`k&}TvmUjr0RQ+o?21+#TAEL;)Tcdwl^AN$v9G+s1h^u(}tp zxpUF@_RaiR^P?MOBYAkdrktH3O!1*0NEy2l=9qi1J2P)cG*vJSG`yN&v!&ORbzY|Z zbD57KwFBU7;eJJ|q^l^?f};7^kQ^?1d;981aRomX3pleKdeu>3aKO6^XC#X>p(KJ!9hP zrEW>NWnR}Sm#0Gu7I;N?SqNONA66>5->_LD4v zJKDA5q3`O&MW()=ufHbg<9Fh1q_#xcm7hXNQtyh7op)96UZ^LvtD4ze_IEq*Ml4Vk=-ply~VD ztE(Hv9LXm}A$$~U0S47f#t>1Sy+4rz#WvKY@Y3cu+#w1U|)vPG?~e8h+S}DvWuXP+eqt6EdLGweR{q zQUKBcU)Dcl(_??n3$JfJ#3moGiuJf}?LMC}fomy+yW%3Dv$#;hpEUpbkMgH`8(lxY zg$+|vff%>udTFh#)`kaye??=R7RufZ!VXhnmKS<+JG91u!;l zvjwNVVUS9(gq-O@QW=#ivr)h|X!qh7x%$mRw#IU9BGIavsCk2ClSI$09W}$~X12`U z+Sh;6v9nV2-y~=$werR`{>$_!KrLbHC*NJNUZxn8`dlyajWfK0h@pTN9*7IZ$Y*PV zcl#*4EF9u5AH2er{Qi>Z6gXiF_*^QYsYbRIE;B#h70yv|Tkj>~=2cx1-X~syaQ>(4 zhmFhU;C@MxPHWu4G_;X**XtJiOKh4Tl0}LFvu{f0%(Hshu*dPPr+l_7J0;t7ru#(p z^=KnhZm%!36q$MH4?l&K3o&5-N=|(Z(`2Rwb*0p3i+7A9RKD0UTg^+4^xR?f6z+A6 zA^1k`Vjo*K?!i4+ku+AZpUiuTt8;9|_O5nE!bRts04-g_s^mRjuk`OltHf9R zBq}|9jrr$egVsL0m1&=`nDj1{up*WMT=1BFF%t9_`{--d=Vw*ji=;qK_YBe~p^GEv z(Si8zVdmyP(=jEe%@)fXndkmqvp;V}c?Z|TgtzU@2c8|z9Fnqy_{53dyrYUw)!F&J z7FIEIr8?6^GdHK*-Vq4PJhk6EmD>7YP4WA3?Fa=3&DPYQn^yHq967b8_F-DoyN;jA ze~tNQRv|GE>L>t%24X^TlDAW#V+r30F5{W-LW1&~p_cSgn==UrmS-ylRYBaTYhAT$WId$%Ycj>M;e8n*C~)1GL9P zrs_%|wcSUGDs^!xoH%gL`^ixN4ci$&5`dutW4`YN-j2zNlfjBp)rovN%Vi=U{e-SN z59*4%L}q6Gq`i^WrbYqLI%00F^?oYMNjH-V$r06fzY$b0`z$zL<^ zsa~;a%G`+%Z?t2+vM@;byBabJ!w^jabt}F{*8Ifo6SE#70~yw*-X;j8L7T>nr>c(F z6wTxbbBmz?T%}5)OSAYk*-}|=3)(>qu z8Mnvwxwa(i9*P*`>c|n62mm{rLCPG4_hMx+>8BtM-+m9rk-E^-E^ysPd(nxgdaNn#WgNx5-Sd3}0 zYxEs|@EoO&q3a%>;S;PXIzs=1_=k6qI!tNlLpyGxf?MC;!1OpdS{IC?w3OWXoWG-c zx3vN}fAvxR$>L!GM0|J9P2K-O1X2f|$k4E1?8^|z>A%^2CT9Oh_GCd=_1cepk7pEY zzC(7oN|{=&>Yu02)_S+7q{Cx8^jt3+@4Ue8Z7^gsx^|fTYP+EC+)^X>7>>+Is_X1Z zML-1B_WI$8Vo9%~iaN?%xUk+vc9RW??`TS+bt6nM#&T&1ZtqO*4^+tOxF6tUhMk*> zkWIZptI8p=W9eK!S$k*HB>JR*mH-84g|4ByZ4BjCq;_?HE@2jc3(%KDv`i!uXBeAx zje^4>j{7LUFCxYIthPV1X0vhZO`Rnf3Wz1l9F6x+^YjAW@OI_S8~T46&8cnNzT-dT zP#(BpI-@0Bf*FnQjo#S?o10pGv0F)L%|BBH$r2fZ2|He$J%Ouuc0#_TY-`omy3*C$ z;U+HAEcUY4$$gM@EUNP5Q=A0#hvfp(qs$9@G4t+BS3QRvUXhCk&nh-$A4@zQf2A$; z=28(Ob*~%HEbs?4f*UDKjI#>4_O#QZftZ~9Ngx*q0r^XuNI2W zZ$(FwWM(-E4U9kb?^tg4D&1@#zPpE1Hqd^^u^N`9=5&?);rCmAV*t zS9_<&48Gx>o=rQu?pZnEocL({j{@XoHR@;$T2-sAwpnfJNQMo?J}R>*8tQYBR|iPm ztT&LfzuKC4nRH}mVr)|0{wuB5Abs_iA}HU`&`k&*I`<;L1|tPGS328DtmWv(kMi${ zqjh$=Cb=&Ju+v_?SY^k~u)m}Ft5U?KqAdU5cda*)=Rl8}&)`r6IQmyfT5u--?kc!eHQ zUZzp*uu(A$(8HN*%j>b|$vrUCW6JquV=V8QO%4Vd*OO|Nrz1-zqAuPA;1%WRG7508 zL`&HTBT6jKTJt>QPN9KnIonmHp5pr#%Iqj08)}N~C<%gx11^UmQ)s9-qs7-wLA2WK z(J9A^xm_J3slo|4;@gJKeML*0RA@bufGzU`Fgeo()nM9@Mw4`%M{TP%F68CkKi1)q5HXG33fe02iBOo==R!MkO~Eaqq}JD-lO}d zcxB)s`nXbR7tVVFDvIY$bxX6r8s3+pfWNL6lYS`J&DvT#Ujr{n=5i9+7MGE-KNxzc z5Nz6KBckT$PTeZHnMy!E>`c0qdAhn4dxq~C!Xw~5mg3K+fEY%*l&*L%?`=xO%oYs>M!LIKI>dqV6<;pWkoY1xMl zs}vaM#xc&{G2Uc+T40;Fwn4A@eNFUDu_@?5(H`O3cpg2ga7ne|+ShG_-sw*yr56g+ z$bwIPvK3II+>LE?!A2-@&@`V_Kr@kvIItfbZw2>|XAJWo3j+MU#5~#dELNcTprG_M z5*pLBA^-61eR({WChYTG+f~zFVf_cZ3;yOWLr+ELlUd2}2a7yu2-VpS3N?2)K@an! zOVZi1aP&AU`rR#85$x|Pa~USSfoDa9=1yp{LiXp7DC&0=pNm-L|AT zS8#S_*8hgz%=~H%itf{ny@q?!kbPNfLqfWFz za3@wK>ogFi9+!C(+K>tMq`{?GyjjVQT%$4<1Smi+xuc>Qf@b%L$#>m9~x(Ps-h;Dr3c3kfZV;cb0G zIeke;(&5YN#Yu$tH?$;YT#jBrLu2Y-(icbYk1|$9!8JjRwvgDeyosUziuj{c;BLcUzMRmZ-B6RS} zdZF)M1CV-yt0-WZ?-1{X;1LQqzX7{qE;_rJT$er``QPe7gT#}7MLBTQ!$XQjO}0>T zm~;(c@TQLvjq%Vh$cB77nxCMB=oMSq9wfU1&VlNR^+m+M1ww3i-L#&<{p%$taNe2* zIa}z2ApC4+cAkz7bQyb~dQIXOla__vL!{ zH-bcM%pl}5ne|e-(G-K=tug{cUg^TU>t|#3`n6MYh96s(d0b82lhzq`B{Pvf&h;wa zf~u1Dvf18dg%TUQST7V6iu1?dG~1o?r_j{U_`~)0U6wq7F)@E5Ppru)T zQ1Oq1_b2HY&g{l?{@i=vvmQn$=zRZ-3Sf-S$YQwoa$bHu^$(aP;Uyr3s(Z*ZjxT1avGc``^xXY_ddV(Y{DU-rW?q;bJ) z*WAq{&sZaq%?X(C6Cj+1S??L}0tfpAx{+S8Qj+r?O)_2p0zekd9}@M)(Bh`v6YV_H zx?|+Fb}bRUp8`y1?E}4<*gBK1b;?$Wcvi@=u64Nhla`iHjVUW}Z0jb;U>UlQ@&bdc zl(`9Cj}qxDCyR_@O;?x=6=+Ks;X|S;PvL?^?u(dDw^ChB$6VR1SRXl=b~0?N8@8nt zch>5*mOd^0=WBCrLbD%9KY$L49v*Bds@?C2oS1m3Ey%pFF5`S2v}{FuA!srVrg-J? zB*-+^MN5`d%0&^6p({WPBsLv<6vCgpcH1$iq8Qb`4<>=L8)G&LeQ_ctt$awVX&snpI{vI=y%J^75 zEz@g|cSs@deNsypSI6Irh;YoJ+@?o98Gm8P(ACYwDa+mT-U_!tZx$j1OY zsemkEke`UDzJ22Lv9m2*l?ij43dan`OWu329|i^Fv1owfFZ-sv`nA<{V@pRoH11=cyN8DWW^@?LaDqeKf2FrQQth~@Vhay5F`dfVnn37l#mjT28jhIARQwhNQZ!Qj5J7hm(n1O zgoM<@9rxb%-u>uS$?Y6kkK0k(=|6Je<_rY z030U0%Sv|F?+kKBz8y&R$I+N3%4b3eOWTn#96HFWgl|8u^ngP|->uucw+a9pgsUroa6$8`zqs=uk{aRd5HSXIOYQW(6N8T&LSAlfHx z$Oc#C(m}5i&4`#N$^EQnZ%*(*vP1t$jf02oLQW>Dr5o2&gBvNOe5@JE9gP6B@=|-) zPl`9yQJk}C5givQxs-1frbn$}JjA78y-2%_Dhk4QoGL+lJlvbQ3DG$gpNvpBYE3l_ zZ|;~0FX+U%NLo_e9Rv*Et=H(}i+6QNw#ag)D8KbXt~=T6xn#X7B(e|a)VG#s9cFK~ zR}2Sw570?6Nmeff9-I0NP9zD$#%l_1NNUGk!JEe8;=;{_H(kBWmd}|UZ!b0z=RrZx zID+4euRK}-4Sff_U#CC~XkbYP{(_hu47kdj{U_x!r^oHwyC3TzDBUPT-3=&?2|;60 zK%P1UfLh~E(waCvp4inS@{mosrWPyN$I` zA)JTvTzSE6zCF{~oR>dqXJ{3vVLC9vmpHx8%S)X8L?`RX9eO68a9{y`_s)IePkJec zgy8vvl8#pErr2&=3yvOoXaEzcyM(jw)2z*}BI%Lj*FEeHLpbB+V`ORMxUw=yCc~E1 z&hdsZz%yP9fZ|Q5E*6BVR~x-wsSVSQ`@C`^tV>!a;8v9{{0^BHNN_p$h=f)rMe-w` zuv7%>IZK#=IMS{L2@x|pYR5{~YK3LQSYpTnk3yKmAaNPej=@mmbzU?Tqu}{+AqH5g z+`s@)SOKq)B_R(cu!0)6Vt}q4-EhxP>r)Kyg(iRAZ}H%9*_rZ-mW_~n3=q_xfxf5v zbYgdpe+^7i=y2}1;woJocM~TyVEYP^XH9g}Qc+30?@JJcGpu%20p-{JNN*C&a`41n zWwKGBp!L}&virlA_fIV%oY}6oUq37z@Yj@OO#m!Xsgxw}x@mhXYrWm9!`xwf3gLbI zZnAYkxP(uP@H~WsCBvu74;D(S$f8q4mr-!btu;`IsX`?*>Q3doByR#QU>WDB(5uu< zn`Rco=t|ORh<%vAo}(`Yfbdd?S;HN@rHE!zkD}*!RJdU)_HPdS^GZm(8fR4^3UsB? zif`3=1wz#WP^xZ|y=u!OTeljlpRZkx}m$I?X}vzK5U z+ww2(t`HqQC-t&Mx2a)(+$5SB{ge%4h}2!I4Y@GDm`XS>TdaER{*Ebxz+bT?Dxqo` zbpue^x|o`GNhhkauUfg?OJF&f_MLTud!()Zf>Jwv%^`Mn9mZrrt)Zz8mt(&U;&j-} z=(b?Z#DutLwRT}tSL1K*t*6tuNx+Krd>EHhyrcy~I;)Y%176Phs!3B~!6)n029%>g z_e0x!*JNH;`N?giY;Iyk+CE_l;4WujZtU|8mG?&Em<|yD4D-KV)Q6vZ=KP)CU$w* zR!BQT1CbC~@YL?)G{Vj2VUzL+iG<^>g{KW!TmIbTmx?8emi$*6%Lg1Mi2PvHrb)2T zlfhT_rJW6lw%4(+W>15WukhB#uR#Ql zZy2M#Ps119O4D9G>88Il);~fIo`zn`eYc*T$r(QxBX=;0rW)N7F=&%{nIzGuIY-FXQsy*6OIA4I647?pZGz7 zf63MyrFc$1jC^IZ5$_s#%mgew z1$fm=P(Q?UUzZH7s+?T)oni^^vSXsaa_7Wyj|aCnnvBM-fM!$c_69Yx7ubaA7|vZF z0gLvJT>FQhpH#_ z@FLU)3k)EmyLh?>l$QGTJxjOsLMl^7>Q=Lg+HwxRKUpxkU8KwS@vmdXOL_FV7YdccoFaj`E4UL0x~y&G|AgKwY5y5G&cQEz`t4(E!PYceyG~< z3U_I$sbmdttY~BqZ_{4sBp+Azgm}0D-=C&uTjokS26!^Tm!kj6)L!g#s7+5#C_kRg z5s&RBeV23{6GFs^!q{L&zuntM%nbv~fnWGPT1M-?LxwyA$Yx&>-*oOa28fT4u`ZbxvaRG1 zH;@;;`r@2o$Ww?v)omjz_rQJs*BUs`QkOFHT`NLekmGn;%T{rAwUqKbUDp(!q$N`= zE5l4^w3KxIH?ujWCkxh{TJHymo~RD#u(HwdkUN4`s0XHPPCv>X%}ci(hyO_PfZ zV@%RsWxay7OERXm!+se1`Elk^Nq*r|tJf!=15A_Bct}y;>2dVRm zy%HA*%J(I57{Fa6W6Ic{d6GO71F*$h2fxCg$MFlLezZx*>-YW4JTU;?vDjvI_-k1# z>ZE>AEO>nfmS-gAD4;uKKYb^!9rab4wCci={B<4sXM2xuU;|0%mDEf=Zm9xy>A$8F zms)j$$wf^|-Rp_8N?(0k=`LB|gwwAt{O=!8%;fEGCBGg513<$~z8_2DKUS&{G5tXt zGZFDGW%5({;ZO!`kS5F?8}>s*7~tuwAd1o5$o(HTWEtIbBU98L%HA-c3TD{>CV_Sd7wryNFDGD!UR~X2PPk^otBMnE-gU#&Qt-4IK8v9>j!6 zP^3#*&Yb_gt6?(lxq4$D{V%Z4p+MM(DYWT1*d}R(0i;i$NCsmJ@UySZ>KR#(aN8*< z-Rk+N@M947hx3`H$0nC!sFbgd1>J<9gP?A($QtMFYM?W14+L1QPhh7zp&5^O`{-_1 z4o|+VWC+gj#WYRvnU7ZExUV;u?5{iP+(*t149plna3I$;8cZP1S(=9HHU~{V z3KR+37>68M5XZG7@FrJg?7bjjkLwnB;~63HI4kiXk1`NK*Sm%6)NP*RtDa@-pDlW$ z%(r6eNPHg7jUO+&k-4={M>Og#>C=0BX~G&IHdRt)o-T1qKOX;DP=hBS{tHL3?|83g zD)xNYid39RXJ$8-HNnF+H0!jPwZ(dmOOS*XdEa0NYJ2smYn&`|ZgFYWyAVxbVhMcB zi5vQtq;%Exa%-jhvySF;mo2w>yPsU4k^ZA5pb0Aldo0?azo0imXAjM*nGoB<25YOxQe`Esww{ZFYXX5Z*7pP~dI?eRg zsl17rWv#kPc~ISzHvec$VHLGxh6m09-Kt)bbGvYD+UZ z{kz;{7(GSp;owzfut4N@TKQiE;h%+JW@ET_iK(Q0Ox2^WxDJ7Rq=&Ln*7RL2IIXP5 ztk*|PK_6^TqblPMz2Mr7_tfmNNoE5fmA4Z8!b`(~TZcvfb|$aV{Pu(0a@f+OqbQlo zib&#QsiKKah(8+VsT|!L3Au~fQ&N<*J~(k)qC4yR%bX;lwq$qL_p>0+9rZoS9%4oP zz##HD%fyx8aD-8lz|eq!SnrWnp3yQ`xzZshf#)!CCnSDMMfae*nyg2gy5L&4iEV|n6u z{usbfuO?r#BzsHg5h4MSFj))s(Zc{tvwBRo7Ee$%*8G-9!xxu> zB*~^Ewj9w5Mxh#U)C%&@@;O;#Pq84$FvsBj+x5gTdnE^$D`_W5*#&>?hK2CabkM75 z1^pBjY#F19En^_R%a~W-Yk4eS`?H|`+g}Q?`V3XXwBve3Xvb;Hb>nX8L|tA!UAXf3 z`>uh;&uh{uIPnCYulR&!ZSMQt7#Sp`oK=kL6z_POC5zMj9IveDrEW^+JBE)4vhHa1 z(N$GNr(?R#=AJGu)}CURQ{z{0(g7iI`;i|ej8@~jxhiFdoQK6VEQ0%TcpK-kH&2!< z(z$Uv;e0F+D%Lwm!iF|_<85lG<@XYz_Y|r9Epg@#$e2SMDoM2#d`>P4^Kbs*m1glhY zhA?VcO~bygvm$oI+8fez*Zl6g9q(w3%HG@p=^sKO<+4mb7gv!HkA3Fp{Y?yzU~8K= z^Ntsil7{rW3yW03za^gCeJaNhCb60*O$e0g+g~)Tx4f*__7L+3uc_cuyXM5PWQhyH z^ke;>c9x3;($2=L|_A%*MNOX~R`f<$Y|@0Itg37Erhp^4gu4R483EvdwZBvt|-V+e7C zJwbiA1pD~{SM*~?B}MMOm&(iUB7-y(G!F8Sf!Gc~c>)GNCk(=tRb@6FkE&MxWOO@P zwfA}cau;7AOq=fteoWy`WU7wZm96`Bo$iGK88(888ZBgQY-S=ch(|(Z^qauWOu+qD z@K&Ph@DU>X|4ssh%ltvgEe#CdA+e4O5=0LhBHn@CE1mnRUo+Z$Vm8*$2}OF+@rQRq z4*j3~)2FBEKG!sUj$ha&&JU6c4c8_umL`^BXgg?1D9!L$ku>79aVQzrG>Ma^duTi} z3l9*#84ohNme0-(=LY^Hc;o|OP|-?Zt^<&S%QE|bdE3k!lVySqjgrd-__P!+%(Rr^ zJ@2G7e6U->1z2w$$nBf#md(g?6U`vYs?Y5=yd}Q1u!zkaMQ+fSNuzrLMI8)Ipae{Y z22PMT!QAw3DG#~U=;DnFG=lL)<-S0|k4Xp~j+bakI#qU+XBvS?r}B6tB!}_hTdnNT ze?@|*l>rjXF$o!-3t6DtjoQTkRTzM0vbcn6qh;zj`eaGtSU-nl*mJ`W;ePM3<<@N^ z^xk@#p6Id?o0}in6_rlI@^ur z0}K$jR-5Q;l2?o1+f9Pp4Bm)*n(pU4=y{=9{Fy;IDs`B#pPO zL5gWyUXiAJNBEh^H%*ng5wwgKThMCs7pRDA#;e}i0P{d z73^UKFRz$k>uBSNbeiV6ma*K#mG8UCNy6PuZrXw~x5$FQcR0h318};O3=5YfB|7x< zl(ZUbp5*uwB5RCOXXc!jhG+=*N84xw(~k%6%a7O^@xGRkB= z>VqoK6d9N@oHyQkCW!p?7MrB~kc(3YyKUbWLrGtnaN!GV4{=}1!vh~&2ZIj19d=Fi z`7P_?oZnhl={ow{VD&llQ71}Q%QD-ZY0fF!{Cth0q9Km+=)|WduiooxPAt1$fpDH4 z>j4qqY_bsQPOul4Z(_t3;Urq@qF7nyc~$!weHbLw9g;fD;g{d)?FT}b8K?3>Y||Xc zqDyWH%?X{;y)&4G*ATHXvrDWu@Cz$*j=fkzgt66lQD#RA=V?5O1k@8)$z?Q^@~Ls` z;1Q7U2Zh7G95@eszxM?$LKQemboIi#59X2J(F{8U>hur;{FqeNi+`x?@!G0Mn4-#7 z?4<>>Hu3ulZW_Dr+xI+O;w#JH^9(GY^7DCW4OxP}lf0)b!Aa^=(w#8Y$TR&Flx{cKLc)#8qll{>3pzrOkjDJ-q zku@g9iD)Rov^4{Q>47D33EaBSF}7)oDaXw5(lruS$Q>ntI#-ZtuCl_wU+9TSww=<}7z!o+T!Xt+dQ&|Pgv2v- z_ge#`RJA`OE%9z?`*LySt(hoJB?f)5yLV|#%V)*4t9cUB|78Q(1?B;6$M6asTKE?p z;{SxMh$gI1r1ULNP8rR*Bk@0VA8Q#R5{lAJFGQdVjeLGVoovG!)&8PxXvn}elHU`h zYRH|PIRa1PR{&+1X#Q@1zg{!jyW6>gR0~SB`QExo=09GTJXxVY4B|l?EtPw8G-33u zT=dev9nO;g9hh%w_=SIW0}-v91>i6w;0qAlBo{9e`AL1DQxhh*E9lmCdISqj8P@XR z_vjBqCxkL<(dY~l`4iq9d8M83M$jSOu$4zPUg;`3f4xKGHGqxC>X`W*j(Mf~Jb~zD zNlyI5_E)1mPFx&%^7fBQn_mJaoTrJ(OYFFD5mD8j6Bo3nr)kULeP;vdobr1XFSJK6e*W z8Dyx|P>{>~?WAxwj_V+N#l#*>Igh)ceF>eBihtd|SS6XJS4weOHIA^?RSG+nUxAX( z+df#GrT=CA^uw=&m-E7RPi=6;pQP6(mgI1z5l22q@?_!Wl}OGdtF?e%R>ixfUh2e< ztC3BIhKm1T7Q$55lxW*TlM z%3!RgB?_-jW-HFW`>EV$<{Ls|Nu6YoCyrS}C#N8@D=BabQ8}zL-)Q3WF#A5@El*a%a zlZT#Qa)#t)#xnYO2~EPk{BQg?Y>&Ro!SUl=6?>UhHub!BpUNTR&A9{HOk3-kG6l-cReEQw|J<#1P0)RKN3=TcVt>PMs>D#(^)R1pEa0mW`tg=xIy7FN>^(2wt9m;GSI#l9Iv38N%67>o?GKn<|8E%Eh z-&+DF*zwE1)-Av+)@c}|f`0v)eNvUUcjy+hz%h5I{E}K7|KuNziQp?;p5Wu(clcUO zOUYLeFZ8G!-}zWH2|jdH`*-Mre}ueX$K!Un`lu@u;zsW51;Rp+7K{9b3GgESOZ72+ z@AJ~fgXnE`78G%w2{cTY$CnJ78Fh3XYOCWGe8DTv%MK4}tMq+DXLIj*S3~Y+UYDE& zF$xS&IfW3`pVChk1BwbYDY8DjUW>OJk{Y#n$rivGM8O*VTkVnH1wnj|D@5a5m_<)5 zu0=;Q|B*^UHD!{B)@}Ris1FZYGQ_Z4(0@^n+3>b&v|6815~@9QsfuGE1D_0qdgwx0TbYrQ z_1t*8?brI9miZazbG+7~yYSuI*%IE^5wjj+$CX?kWh+eVa8+6)5R3=f^L^$RAf`}> z((&tk4>NznI6UBv8a(p>6G~zOY6X=~liw zXXbUD*&Dbfz2cfNrg5%YVxqb~Hhhe_eik-Pjy??8!f!Y6EOK`&nu2F);YJUfklQsJ3 z)EFRE3IqJ2!j1sP>I(4wPi6QYbN0WL<^M0&gsMDp_;BTyafQ!izJQjK6cimJXZ{_^ z4OU)nDk+q}Ef+dw5fAVSMfA?yV^9)abKc^XMIxNH6u<<8h1*HlWm6=^=5*S3+r=RJR9k>6!I&InwVj2^tnnFm>yb~T`-WqDCPksUHm^_;jssyIv5BW-q8 zg8n=-vuiPeZy!No#HhldyXZ#R<2W;DniNbQ;{Yaf-FFxeW4VB<&tr5hJ#3j7K#Eom5AN)M* zIGHtiLB$hM8-FF+C#kl*aDxgu@MV=U)qRsE|IPP_#-c)oy^*g>iWuMp?zcCYs86Pq z-tVP0&WKJ*cF1IAA6FaMD86PQs_>4qD2HZK?IN z?n~}ama`)^EcS@>M)M?J8h-mkQ5>m4(Ivx`;#WqgQT(fDhRMMDApMKBP&dBoAw zSk;>rm*vr&NiRa!uH&0Vs&+5LNg%xEApRK#gus^t0RwQ2@}@LqI9mL7QeVclsW zWvr<1r|@kgF;*i%*yqc{tqk3h(wd5#a*Ywt)L>uExd zSrgn&Joc=$9%ejSU_P(zPU1tqnm2}>(PyxH00>+h?nY23R>wm3wo-=#gyP^DHN<`b*qK2I-fJpjimdLcq z-9cF8oUfF>RCi%}$b#?5&+~9O&En!Zdi*b zy7x!#4Q$JCZxJOq=x)2V93;$vBSmExVDUY8w#)Lb zmu_xa^*O_D*E?g$-mr~R#Kof3EHVp?18PAXdph=j<>%$ zs8sQe*KaJS9J{NiQL*kp!n%VrOBgF%)jN$35D z(pLAc8&J5j2s@a)Q*PoR@sIZxZZB|53b@&CiePK%osorJr&~sSqYrVFwS&zchGw;H z+LwK7D7G>A$`hRz4gASs|6ji{=3g)CS>*vad}cZ7Y7-BVzsZVi_Wvis4uNfY*i84J z6q~*MpJ9OSh$m30?TF&JAvEK36nwTVPGL2W>N%$wib2;B=h@EvbE+LA9MYBO1vN z8gjB#W>FI*z@eYbwt5v$h*NH2^_XV*E&;tc&~r0=J+>!Ej^3SI@xccC>`z{&|KJ58 z9Q;q!$4XzU95mNQ-KXod$;Pu?u7_=<{NT=~hwR1X^+U#$#H4gy5)uWABEe}o>||8% zUzwONv9aR90P@_iKSGzCE6iPS${}BkY^SZ4e;sCinBA{DNUzwP{FU!@*-wAHziS-) zL*<+JqjbKbJl^_4sCJe#@oQhM$q~i=8Q%XD)sBw-FMa-9ozJ~q8#GTIOBV5hq&pRA zw_a1nP|NZL5lbnnX_z2Hdg)>nO+fkb9!+m(u=t=tAxDL<#yexfxpjUin6teVK@0NM z%=uxC_HY8lvzYTOt$%aJkuY{)y7t0Y(u84FzuOb3@Tqp=Lc@R0UE`tQ0?ODuCnHm= zkG=fuW1BYEJh9hs^HdedK?^`< ztZ<->J6>qOB#4{d+y!>Z)da}=1q3O?Mice-+~ze8r`ilkMZK5I-?nW*=6P0o7}wQl znGK8y@X~q5?6N^yLl--;0*>HM3!hzIvT3mObT4$h!a3Ke|6`=b`~K$}2T#s0tD7%j zd^MK|0wGTwD(cv}e>(TDEC9bs(=!_?t0!0w1s$^LLfIPYa>(nL$SGa{?FX`l)$|VU zd3v3^ov9;OF7ll2gi&6w4gtH79;MOQpFh`B)u{OBv-1ZLcTtJQ@)Z2S3)CF>y{JU! z(O@Zf?DL-4d*wl-QLkU{7os(;0-|Oiza2Q7MUsWDp-X*RCxAzhm_3Lmtt;lE$&7-n zYrdhX12HbQWw60DoB(Rd!0wwKnWPv1E-QHhY*7AGkfrmWfDRx9Zrl9s-s8Qyq?$!4 z!%|jO6f;z6`7GgYza;VzFj13cfSPs`6t~GYEMA;gkA0@lz{^#B+)-gLBr4x<2nSc6UxN_9J*P=!DCALm^~-_3*o25nDH7MK9jC_{pHV&dmdPcf3|-YFByI znV&sb387g;uVrrvM(;n5*K!gqIjgLWU(m71^nk4im+9isY0NfIqDIpX?@=j6<8ePF zr+XSx@WYlfg!k=$>qBf~gu_IRo`@ZMy(uel@rOBIyfx)g@QAF8WpYk9rSBQ^uAUmU z&ff7*zms@dN`dsQ=tN{|LL%vTXrZPuzH_Ydl3w1VM@5xa=`DL@2FZ9sw>v=`5I}`u zWeI2aDK{QTTMxgJdz#ejj1nI$v1NRpSU5e2Vy?OH0A)aztki+ew2ID;ANL(Z&ZHuE z^&PT$odoFG1k|(w7mEMbNh9L@&dYXtOklwL7sopHCFoWpz8&^ z{5~6)ql$m;YLS#iAI;>%{gC$CJ?zr0s|;lVaUX~Mlg74w#hY?oofF})G*{>i;&SEki%Tzkg$oi9$7c&FIxYVq7TMHXZ540^-k@xZct7 z&ECPDGv9IAhl_AwRPno1o-#=}crZuhy=39sB`l75K69#ikbCj)R}$r{8rk67N*||J z_MCoRePuND^+YLBT%g@~ahu`EhoW1YdHFM@w{B+7wN(HTC`Sa1 z9@P_P=|ojc^AJj85^h~5iDDkofn5{fJ?shAcb;STBVe@&)P=Bva7}DgC&wvWwI-cK G@IL@OCo0_l diff --git a/images/custom-only-performance.jpg b/images/custom-only-performance.jpg index e03364b1bdc31b6ba2541ea825eeb6cd8792c95f..ac79a3b70035c9540e46d0059b625ed6d6e29b89 100644 GIT binary patch literal 16755 zcmeIZ1yo&2wl2JI0tA9X&;SYU5_}`UgS!NR26qU)2?P?{0t5*pxLeRbu;A{2jk~+s z-v7?&{`;Kd-qZJWzt{cl81F7dF;=Zrd###Pv%Z@1n{&Y?VGF=Ld1*Om009925a53R z3<{7aifcR-m!Ot4wl}jir`6N?lNABTYAJ{2X!eR6VYIyQQ08Wvh|a)!r@ zEF7FXJUmqN0wVle!ff0;Tt6OyaOch)EKDqtyLU;r9*{ra`iBo#Gk}YZkcG&Bgg^rz z;vyj7BEZ@JD!70s2)`Y`Uk?OCBxDp+H1s|_Q{$i&0T_n2Qm zQ1XeCw2Z8ryt;;_mbQ+ro~fC+#S2R-YZq5H_m>`?UV*QJg5SIi35|_=7oU*$J}EgP zGb=kMH!r`S;&WwHbxm#Em)37>?H!$6-900tW8)LwC#N8bOUo;(YwH`ETZc!-C#PrU z7nfH*P>g0>(Ed7$Es53I@Cy z1|MpGTg?|8u1I;nH`rk?AoSziFU7Q6} zRytdJg4sU0%+G=rH;USddKY>Gmq{VvfW8?`#B4+a^b0K*?j zu4iuc>O{Jb5MxCxcOKF{K0(X^6c7JGvxng$V(9M$!)bGq^0P}mn#u9>H-x8-$ab_y z7niw5h3{0wPzZi61SkC>JiXelC{jh*j#na&dPZ`u*DK(CHM#QJ!ms zK$HgcuJgwcI9DApVhE-_FrYwSb4&$wTz$OqWW3U5mrurWWN%KEbj&QzjXCCW9H+`7 z7wUfng5EWU0V`gj<4fsRNE>h{)9-yn*urLnn-&GMT%Ej4(YJ zb-|aEzhUwJ{zJAttfeI;>iO*K z`2NQq^8QP8xkJ9y86+XeltBA459Y(F>zAq!!zKU6h4_*%Ai6W8X@@TT zO3p#wE#wmn=rV1aKf_bLrK^E=?QfIxpX~l^2e~(|)?q+dJPd%w_rrg^sCxIyr>Z~uoVAM1m!jcWPUDyrJq3;% zdB)CV&8Ec%6)+&}emWGB4SX?cJ*<&xS8xMfxHq<+c>`KVGCGK!0uSC*h5;yvpd0(B z3t8Ky_266>5YU?reXLih-XH{Of&#^&+k5b9=HOOIxNZsaf^LWmV8G#*8Y}SK6Z?bg z)*yaS--hV)X2EbMp><$a`=ez$;?wiHj7HiQ;1+&g%+87}DETP6d4+88_EUN~a|~;oicU5R0n*_V%^XyG9A0iM9oDr61PV5b zK?-LIAMk`c%;HU!_W1o?@ZrU6`@NXE|_8qjsxtM-^2Lz2Z(3&-| zNk1Yv7`%O*4g=)a7?AACc|ey~9x2dPIT(oIv6mb;00d;p$x#E zm7j5MyhbOD0It}=FyIWY3I6lCA9m93$n{4Q`-jfUS-~rwl zRx2wv*RP+paW{^~m(P&^o$QV)YS*p1;v9wU@DQ%XS%4o;;Fm5X-j*KkB4oj-`S zY&sz>gPUsaV8H!fn=>naxo%pBJh_G|B*F){9f=CJ4aL22Ztf@_=7TC}8Y*hzvA-1x5UrqQe?ryE>(*T{9zt zQkYPGRf=%Z`JZB2i9V3*SNB3n@0+nDwMN5yd5LvJP|ffzI(!y?sycyJ&=8%zx^9 z%rlRvr`u`G_O>CRJ#?y{Ci+k+H3L7cHgEZ*BdX_$|MLs|PCT1v9Rr0T4|kc`YGORA z3FpLm_0@d_e9wn(oFY5&tkQ9-gti6t)++->&|R~}W>>3IEb$yB>-JCw8_@a(fOgkL zilNh6j84T-vk9KwFiZx#8K%ux8RiTI)#J+*q~`AF3(W@$EFSLLa?lYNVDt&HN$C)C z?i0hl&sm>(+4NG_G=C5v=4HRXIWh5gq=h!nfxIuS-;Ou73x(SesQVI%b|xdMk*=Yl@w-`L4X-~@MJh!s2Qz8XRk*YQFgnz0t8PGHMaJxWIdakbnlu6`b>8%cl_>#fbaVE%x& z)A&t_irge+-&yq@Sc?~DNp?g~yh-ZTw^Lh+Kr<=m4|c|HHV2;^9!Z)mg&L1qU!NmD zA(yKvMt(xcQC2s~&i*NQb-AWqpMNOldz$oAjis&Q(?I0qNG~;bA2s|v(Uj}%Cqksk z%4bq~ZI8L@-Dm3&1O{#0mhu!cwrD7{EO%>qZ>tDza2-_FrNiz~lP};`H4g|N3A?C* z&9~AIqN#B2@~O<*Qf)}<+>D~X?yMK8wx4l{n0NI$6r&K0&zC8f^8G{``$(6PktU8T z)6JwP7dbLBkDnGb??N3ab_@e17pJtaH%EC6!_tLsw=^O^be*m0PQjO#y#pnif8}w zasqVV@)8Dg+SvcU*C?y{VB&*~CkmwMSCl_+IzvO&9`9*`PS;+tX9abfAOb8XddoZKnat&SRG?Ix4a^@xYTOW?VVEmB=V zg?r));Xx$l1%|{GRipbpGB)ORcdNcwD^xI;qKBpEJ3FEw#VRkWRK;1uj@cHoduXCZ!d_Iq4^#tEcMcbuVfMKbIR{ z;cB}6WF3_vGdnntGoG7@1&aEcRQzN*on#af1URq7SPkFdITJvDF_!W{UDNBa9fQtqrraL=Y0NAcF20?VW4 z0g||Q-1*Ow=(v^x2c%x!W65{8KSoGePaytbnXK z5S<-D7o-BL#L5ELSrZb;AGtR$-sE{k=5V>|o^(~S$Q^l=xvGR#G^Mb;6W5Vcv22@5<6GD{j;y7 zWnnk`)7%gMS~4Mxt->HjNcge1eKFac=V~J?0aBdt0D1L6Es20gp<&%1MeoW5x5~SI zas;1>7jfxyo?%G@X6E1b8PFd3BZ)8|p?%94d8V)AZh^IDZu!`@J?n$Lu~@13GvHO0~zu>&2iF3gZvddR7X$J_2m3T5iXR_ zZ56y{m;HJ30Xci>#5cb3{9*1bqZQYuC28ftaA%#`69()&1jB%A7Zp^vUjM`|{E=(; zKXz_0<#(<#mb-PY(?0ihdnEZGvEo$0&N$V;ig<|b6<>-?ZH>fwiF2*F$mdwj8*bRO2hGTRh3dwp>IR#X>t06Swe$|n zF{zmf4hFKZ?D{^ouebbf#L<6;T?8?w(0Rg3YPHp@!Lgo!ViM`q|AEb5W^a4h zsWfC?f|B*p$ekOjxSES10B;=)9my=!6}5^Pvx9GB-b3$dWbSd=3k?_`BlpC@e6Bva z+#FJbpyo1^h9iLxvsj9~btPKzxiM-{LFlEHZq{fPSKuR#UGj z@A*Z}Cw{fkax2|;>zA+}rqf!U^f!RUvajm#-rKr?JHPieMtUo5PH!dtEARh5Srm0f zB%AlEuB5L-(=!aZsbx!)KA~G!z7Cxhrw(Yq6F=!%Cmh#uuP@+>%2M#AXoApbMe;Ar zmvPq@u<4ieVx_hE3;5!pH58K0&_lq@0O4wm#rEiwu&FEZ zCT);Dk12+xXG__-d|mh&D??6h{8S!?e~;P3AJLJ^D!?*a)G}VU=!M^6)Ii-nIca1d zfohTYqpvl=#*7d6ctX*Az<2nOUm{1I4zwWCwy9WtPe?G7u=Y&f878PsleF&h@91mM zF{OJEOK|!TM+%t(`JVU+u1@3697X{ZIs#?7FP6AekzhoddyC&z2l;_WpgrK;A*3OS z|2b94&DBk)`t{c11>N|t9;1_dY4NG1`Goz=cEI&SRSF>ZmlhQxgD@e?l+18Rd|8*B zJJ5G26}2OFQ-rP#X>&`^-b;C&6nYsPK+xj_QMhY6G6W;YB2={aU@-SR(h+dqUuh z=`?qyQVM>gY*jR=KQQ0FG2l9ltTs-IVO@L0hI@?F;%Bb!zpYh9<`o#+(^TL{Vo$zZ%Zh?#=?L*9a zp%$jKL9|TV;b;Bs?-0ljBwdhNobVK|0&O~Jg0(1LQ)U!I5X%k=%u@!N?P4D8`v zpQKiV#aueK-pE2r>tzkqdCv|EfK(BRlHFIk7~L(M&CmYLc{i~HACVkBzI`1BCjnWv zf8hX6eu`fFHiriLVGYdYOZo)IMVWQ=uL8v=R~L*LvI2T=kl8<-b?R*t(zYG9J1kx& zB-~k_8SaP?cO2sG+A9d`Q^N>ocXnwWG<=89y-7*==Gz>@NLI=>FW>xHPnysz3N`OQ zvP-2&j15%zLYK`nPvDBe>0nbUnf>jF)uYj@m$h2K%ptr2dx{NfeC{e0v3&2O$vSXM z0N%;J#8-YL9O2yHzu$1Nm`tzQN=$5OA-!fXyjFs9)$=kAH;B1od&z^!2^uZ${bLxR zj3kRPYf_>eS{GXd5T}p-o2nc(& zcs%-Rs9iLsgVai`@#b39>;xE5=fnsE?|5A8V zjG177$UM*c_!)l^0~k!dWj6@Gxiq&-^M(P9QyNqWaQeq*BL0QtaCxOookmej zYo9g!eYi7N9*E%;5|Q;#P>)j#AM|@Ew&Q#@y_+T(80k!GMea6ey=fCM>_b+#StH2N zaL7Jw>-y;}Yh!CU4CsCI)@Lzf;jzzmb1_oG2NUX&BYGP<#}sX^}a_rv7_>I#KY(cq-_QsYZ?2x-d zjLexVS+Tl0_cV^Ya9e1z0ri8TimHg+hFCcoJtJiY{GRsr9mPihopIs2-RqIorLULA z8Vd98?+r~-NW%bI#HrUGZwm}7Jl_kiU!a{9Uz3s9@eQ-u(H<>I$fAVd$e`Q1{FzB_imHlsS}a)w6^NQ^E>ikbQ=V1UBd-oET5Ey(=MM(#Q& z%n1eng!tcQ_pFp<`lpXKtPozk`u~YhT!Z3o`m)s(?F1`6Q7jfxXtR=xp!Y=}tcNo% z=V2LK7ozS3hcS?BfoJr7{AtUSKpOHQPrXbC)-qT;} z7|9<-bvKW?skm_~2%#qttx*rs^ja-#~Ts=+mHk z@O9}w2i{K}$~Db?}PVcQBaboDjmOfO z@f9DL)EcCFeg#9TlD537tQZJ8%NwXZr5%W&`}T zpZ;ni{j+t&*e9X%vC@#Yn4BX1xmxx`ej^V?HI7{;-(e3$8i{rW!BtqJd&(eBCRvLj zvt!=OCZC|_yZwvjUjBX?ASnzGk{*Go5>rnLql;WwMkr$}VVnb8!C39}GyA`WaD3Ys zNk5Ia@EyF~%bs92tE$u~etT^|Ndj-oCm_YVy+By=6bX4kP?I#0}pEej4fti{L-y7YLQrgSWhiB1aDNWOZLE zS8_h-nhifpHA8epGq}HE$4eo9v_*#LtF(HuUB9N_D;_WWJ>&t)HTDc47Zz$KE&m{@ zzfi$1+m6z~_-gxj{Mci4o};Nb1)Fsapo_3kO>CK;njE2vSEkk3+E8=M&dQPnH4|Lz z0{0tbHESVE&(#lAnB$Dx z5#NVlYaI&e+4`IaSt%)#uu;cbO!D!WmI?0?Snva8ag%@`fTHsAuJbQSu=t&j{!1Q^w7C>Z8`G;;xn{(Q;?Zlv zHRel#+eq3k6MO~^4wqXp&;StS@|_1zNS!4?V6Mcp#dGlYS||dvYk$k*wAhh@gBcWY zc^?Mc^=$eEMOT#~{1mM!^M6x|Cid9l+?xdy(U@0EJf>>`M+Fh;wxyA0tes|f*s zdwLF;73J2JQKeB&ZLA3guw9OJkaw~_7HcDsUm2evp`@H_Z@Pf%IRN}1NaYo z>>pwF-wUmTNsscr4gT~>`9 za70_yEftw4$`O)^h4gpa-|YAfPd0cucEkmGyF5u*#UrS8n{8ie{`e(h4-tth446^! zzdrKamWa=0J&~H@n1&~-7&bYPb96c&N@5^K@?sPS0!;D@gW@wE-ij$xCLT z4_DKqG0>K5SgVAJm0blj@SnV1EX@tGW;m6yH@I@DX;hX}4fHypQCA(!=Y?id*WUL_ z&1b#Fs@|_dOkei#3mdz}is4Jo>3m+ncb2?nF)Q-zc9ZN1VF;m_;2IO1XWD@V=%oAH zwa!jUS=3nvS`V#K?AA{Z!r*+lmH^M@_-z({EWQ8MGa25K$s$IXHH0{{{?Sr~Wmf|qn+m&9+C4W4IIqE*3hbGCe%cNVi}nXh2(k6d=eATm&I$%3xaZMOJx#4yks?2zLi=3PHgW@|**WLX>oBWL=8pyw zgcJB#v`kLp0#GPhN9b=Q<}d4qey=xERLW)WISh=&1grAFa8! zh}6SMS!N(J-FD2<*_u6kNPt|2G&Yx`+*$=szx1(8z}mnwB+b;Y?K&+TJg)ExW!ijF zM=y?D{bauPu7!_7*6;^;28EAaPg2ZQIM(qa`J*ZW&R*lE=C-@k3kwV9S)<-u86&d> zE$~2o!ZI}Ig9G&T8=S#m_@&P57n-6_XD|dcYHR-G)_>BJm2I&fXZ7YO&V`&kXyZo` z$o>rI#yRTeYRyK*$R7ERN3Ocrd8^kAmwp^Qxobf-#k!*4sY`=E$0SkyU{Eq}Xnc4{ zD3~AqS_i!wHCz}_#cVH!_i&cBfW43$!;L7>aqW#MJ`@i#{fd}3>Uxa|Dj;)1M@)6X zwoP^3W_06Ry^)#l>{~LZ_ws%E^f}cDyUky&Ctmx3uZT~<*GueB0eubiRbLpuencz? z)x3ZKNny85pj*tVrrSee7(kx+b6eG(>!xovfgiw3H{iZC2i_0z`o$unMmXQqBf{z9iEU6hR4|(X<)!|e3_Xzb@-Pzv=3amZy3j%8}|?3 zrqH@kPZQNKXH(4@JU*WbXu<)S&kfa=|TB; z7R$VMtr^w9{rB4tsAD5ut?+rnTd>}So@D|jt?Vmoj`m<|FBjz;=1cPKi5tZM{BST5f1m4bqe*{q#^aPR~4p04yW;VTz z_Xi6FKJz8e6KNubnye+tT(BSUf6Mb|NOQG008J(gfgkhY8+w$5f56u@-8|bSgi6$h zNVTMjITIn{Y(yWHo@f4d*tpv>mAKTgwfgE9b; z&n=j&8i?|}n;5{JaHimiH>W!{+~!R9K>v|TI5}nfoSHpZz=}OB^@~(+oJ?i9J7o_= zOdxjQOS^laIA~kt(t=DP$*(T%Y$@;d%Rqw^O`%W-`0xrcQ(wJe-{F+zq@ZF|9&%fk z2Lp!ki@@{K%21HZ*;*zHkjR1aHefr@#jV}27)knca?~wGa_8-BC0un+!1mUl6DIf) zIva%nQ4_NS2){f-W1rW~`Nwi-;4SPV&_>zQ=(iN^y2&SCxP<}dQ#cOw7uw2lbxJ!N zmmb0=`(h+bfAz4m!zSraS2q5N!e%u3Ng;aJeI^=s zq@F3S8EHrbDen)$9Y3>pQpknKdAw(K9LkSz8oe8S7^nSvx^B^~dzaD1Ifw9ZTM$W2>@}u8PLHw*j?eS! zi3nN-H>PGnT4dZ70)&>0v+A1dMTKYbB)L_TRL8WTZdgHw=%#@!50pd#snnmL1)3!& zqixzT6N$<~+Eyyeb_Dj1mXn8ShfjLzWsBKdx8;gI0a94eg){&MVgLvKa^$#^rm93i zkO~Jt0;nJk|25}An|Yb}0pkc81fE&omOmBrT9mds$+-kB=c;!QUkbWD8)ZwOKOd>xGCD(hzO2TsXXQw^vtX6uEeYik`{i$l!?KQ5F+Ad zKTwf#F!0LB{kzMCmu(HpG~`8PL~2dG33s~J`0s)JI#Rp#H*KJ5hc91Z6(2201l8Ol z(kXk}fN;|-9QVvk`uXb_)|OCi!B66B!?ldmF=@4QNKwXO^9IJj?Q=<`dCl2+R2x!=YOFx~5}v}eK=8#FLW6WKiZTO!cZy#}Sw zT~c#u!DZ$>?wI6btKzglsfnGBrYve6p(GYI`J{o&ETqRcYFW|f%hWN=r3wYRrWS_= zPbQ*yUeZ-=m?F!|0tr;@%}8m@DyYOk(Ry>+f}S;%kveSmx*xu?Bj6&6kZm)6Pm2)l z@40YZl2AA`39&lxB+@FImW-V<;CyY6mW%XN=WnBd^`{E*9haggIU#-FTF8UtMG}uB zdtzS620CY39fklKw!lu&T*g<5cY^WHF=7JH7ece0GOb=FltE@Tt6lV)`PvDhnORIp zT7Of7F$9H{BbG*-5qX{>qA7gWTo(C=to0@(oE&ER6mly98*F< z^FL!Z|9HK`*hllhVaQ1Tm%IFRmHI`gDIWtQ(%w%hFJ)=VI37vau|&68#H&<1K}YE# z$4W!>axuX7pi-kG(0(Srvu!;Op+TbY1lbP0mwCGEt}XP%TLn$sZK;@O;A&~Wd#Hw} zXHf9%Y(uurO{>nUS4Ym=-&4&(OML7WZ?hB5IK3En0SY@w+SyARxZ-4SuP|skztEd$cb5e_&uyV*07I+I3_>BbZnhlkm-GwR}BGnnsBmz!)o z#yzi&_*$rMQT3<>>&u%r4tusIr>dp|`1%A^i2fG7Bl$A|7@VDkx&GLp*KH8d6px9T zmnjd$a-K@>lp_swH4}<~Av8r>SE+i^F)n#M*(2z^$c81}4CnNKuTAKf4=Df9r;?zu zu~M3Gx?_`Tkt@l2v-dyAy?QRqcd9alS|hA(O6P~-@jzm&Prpx9&ZHOX?GBo=1NKZE9&%*Wk?tNA6h6h($0z z+A5A@P!~%a#mTPytmkpwsX=Ic&enupd;Op>X9ArIsc-nlXoO{Iv*(f0;-@7< z``;jGx#sSm3O9y5VuddO+`f5P{=IbG64HckTcF2w8~i8xv{j~=#h*mg$()F(THTjV zK4NJ3F%FmRCyYQ&giCa!eQGq`m8wKafA7GVrO#G77<`+1h@Dvwf0WQO;({!4d@vUw8aq*#rGVM*VIGm`ghxoJgQd^fIj7l7U+TM$tz$5G zSAUn&)6XxP=kTU^Pr3sKE%E~n4}zr1_o}Y~WU960RZ&X%?bSM@)4SLa)Zwz|m(FrC z#1uknRVTYj3K8mRqa}nyOT*)1Ibt3LcoBm>!RmW;%lswA;F`S~+3%3jU{UOk*9ndi{KEAvAJS5@aYc6Smx8H@d zqZUc|DNq%4{$Pmw;3yemBnc2+t7aL$uSZegYFmlhQzTINySCV`%;t|eZ0i35p~GQx delta 14780 zcmb`ubySqy*EfC-apQoH8as_ho9Q`R-H~s_`}m~MOu;D3vW3U4XW$4AM|=&iB#x#cR-lg}zpvbyqBXkI{IN7G zw~xR|D)}z;LP@L$!-c+I47s=G=#si$8mJy(^P{AP{g++8ik70(?I5ul##_PYH3FEA z(CDBB(oxYx&ZklGr*pe^y8H)x43v6zVpWGl(X3&gg~-6BHWTFu_SW(F$M*Lb3T)~R zQ{Gt#R?bKi;(A?IoT)w0b04N>Cqsj=mlrtarMa zDY;`To=j0<-YkU*df@>yL+hHQZ}^B&6x-Cva^AYI(-dwbAYJTYVA3 zP0jHuqa+Url}*>|;?bCV7c74*=skh;GjoGvzs5#2tfzJOKQA;S10NQ`Xv8`858tgB za%KnN{vHJ-wi_ijlT!}SBYCr}u_A?iTAPNfM2C%0k@)#5QS?6+6uJ#nrdTXY9pmDf6V5pt9vgPJ`4U##;>`e z2jjvOhEd=OKgBz=-MCxhSVZ}B-((1QkYlV7B(QHUfCN~--=2kW z<2bNB-<<0Dlvglq2epdSelqx-H*s>0hX*_Tv3}NLA!>TRa9|E5#U=#PtlF1sO zQ3rEbOKJ(L?UTCqM;|e7~8ubB)wU}W-RB%VKE6X zf2ed-f?q|He+LO{*F_@OSDd@j$ZB67wJ&5`B-d&yIk>-t9WXP2&!MQ59u*^jAL6Jr zE6*Ct{`(WHTUSF5@hJ6^W%7B2Wh2zA8N-w$uk9MH*~HC{I#EXWY&DDoXv?77u&6ro zm}_*Jv}5qyH7O5*gjV2Uu?Pt)RIed{D3l#=2$HV%$5Hk$c1Hr;TL$4?;JX)RNT7qf zU^ZZW|53%cszdu)NC6TE9>_pQ80-iveu|nSYt$;h8g7UuJ;&0{_)@0jLzO$@n2q?f zn8&k;1c=T#h_bv-&CENKt+wmcoC`Hw?a8~Up5J3T=VD3U)C-pTOs`l*VA6f*@wRN5 zN1e8duA8fz&+3lq$m_=2AjWDp(c?zRAwvfAFh_sEqQE})Y!zqx_)=4|o;y?qrUyry zPu)C!MZ4Jx=$tY0|9HB0Ilo>>67~OXxn4Ir*IDzILeL-*E@FI^z zTEAhEq2eKzRh+V9j#|98M-R6}oQCHPzkN&NCL-{b@`rekb*p3Df!-vGI2V)(OZH!n z%BykUO42GyvAyiT()};TORG{`vyxL`5APP##e2^3LUfC zu31IqZEsvI#0@FkH8>LEzF?V>s0%T1cF4lkZ>weE)rBo;USOc$@Qeou-0&n<*g{jG zb-}YrbaB^vq&;^fSaZm%B-vd0l@(d4{5eW&%gEu$&2O{XHI%RG?i7CX|1Lkl)%R#E z3#-5v*nlZ2(&>JKfATY8LIUZNf)xvNLN{b1g|0;#d|NIX7hh4ZvdV2|urA5#y!pxC zeyqoHJ)^VPs#X9wbEZ5oFyAs7C^<#~BH1*YIT?a+SgIN=EY$6LT=*3N)}x5UfO(P>aP+g2Zk;inkZH5-k<&8=<7K0Bnl)~3kL zjPc|wF_o_v1g|k<5F`YYh~QJkbCQf@m49BDXM1Mubxk~%@dxfA?Tbt!;5sPq3<~Mf5nd9a?ruXAXQSU`_Rn%YQ|15Yb(cj zm#dq@%C5VA6k7~BHr`?B&+>cS^9*auT3^6FS}bv0$tUR`oWX@oNVN0o_Gt_}ld&HH<9UgUI6yTGt)EIuQi z_mYMi9$uJ73}NpX#g`=w>2Eq`cK&TE=+11FyFx!E2t<~|3j$uq&78(U6$#%br?l3O zIihP^QOb4vJPW#L>YyqIlRBFj+Ov!{njMn6Ke|4%LN|G9Cak*hQVCZclnwHxmH}~d zWpBk0&e}qz%(sqZ%vgsG<@e&*Z&oGTGbDD`B0Z!Q&{y-8&7`NgFOAq-9G5@F`;_G! z(&sF8Je7C2{K&)U@hboO1m~k0TsbKgx^XT}gmT)^gwNY+n>IJ!am0%Ahw{(ex&pHnRG8Oa0bM-C!&V;e0AhyOyRI}alkb~)I zSv$K3>G^YgC(CHH?#wSLQMf!3Klz34#j7&>gg}9H9@6_rAhsv0eVZ-gQq37UDItF9 zIcjRANp?qB_^^LbU!{oSmC+N*wz_cve-#facY>joP4((KMwWz85hgA!Dp;{9iwgAQ zsEkAj(m7@2+Ub;+<*7U$HQAIHnc-WALswzkJqy9yoe>uo#Fl!ZA~`e8(QfZjwILwO zC!!?FVrfjndxxHIK}vEV)3+=R&xHh-d{Ehu9Yir0cU0c<)sO**pZ>)WV z1X{eDFZ**YN35=fEcBADqR1HIosq!Qd;3;fl7=$|2_*2zO!K(Z{4>YqOMxXt-$2Oz zY808a3Ed)WgpUa{PMTw^92b!PDby0Z&|O7Y-vQ9(-H6P25F zdQeO(-Z0OZ)pO?w0xYt;qrjS#SIBEVR9O$Z8sQ0is&|s@Whi9#?Xwb@?j3Cx6zRC$W!q?ZnPD7ki>I)yb8CyD{G1pz0eYX?t0Bs#l%Sr>ILYcaV=ye}9RqHuW7 z_{+Vcn-jySO7>sKPI!B0M-kypXGqLLB0h6(MDlh|aV4gKGc1;j!>RPfwco~Y+S~JO z960yv)ny!%Jf-w{ANo1tWYSe&oCk~bP*+=b(fa)-Q5)q5I=hH7JwQJAoz{Ac!1qO8rQ(p~+!0J| zPOZYgJo;Qlgs)s-*N_MWqqErGWWArVqxNBnGqLACz4&~~=uTIN3MOIOEx>)!_f0b6 zJ52k6+4))HT){Wdb1H;aDudWM52JH~Twl+Jcbb^+l-QJ4WWjd};?Fg`WW|#xN@G@p z`NyZ&4m)H5P2*Bw@Rlh~45BY3-pzW!>$5(fZPIqTJ>MF(gn!Bzj{ncCv!-s5Aw8HR z2|apA0ZbE{4wf!J@Z5N^y%jFn+um%Mfw(YhztX9-7Y#i6S_)Zv`vJw?qrvlu2dC9_%I=;vs=JbH;og<(k@0|a z)2MQ=bP%Sk$b_mln8d?Zr{)f))cp!s(s6DxaA^rw9y9IR<9#_bXb#v5-Ek<(s&CLx z^HjorM(l3y^cYxhQ@gg$hWYU^gNA73>Up1+8WNZ$_f~^Fnv$o!3~JzeF{nQ9WDh+3 zk_-G;Ad&OCJQ~>1fsb`hjZ4V$$J}F67(mAcePyezdW3J&B7t`pwo~Dc6>kSgSKz78 z1u`@oX1B-q4X{St6NDw%Qh!WS!QfR}t|G57zGOQv-MFJpSaefF>(sT@hkrHvA!Y&1 zd!u7#3kq`(fPlvJSc4N{YG4cZJ^2Wlh9fr!?r7;O=~_*)JKX+b#{S$rM3A2PlKLg0 zPBc0eb`jW~m8V3e|M6q9Gv#HquNe(XF?ST{Iq|VO&fp& zKIWN6{Ay<-(=(V9oSn|=YK>qM#xE)AWyveC>$aWOgSjv$z+AZ)K~&vGrmb-6zYYA4 zG)8PuR^A`q7sd%SsnIJdQZp>T31!V}nWJvU6n+7I@>Z+ab6jI{H^au_8rFHT%B7Kj z=>%if!H?-DCezUS-ig)sSKcL>o$heP#s}(Q-U0|8HpmrXxT`n7--Ik1ermAhs)hdS zs%Mu1H;t9llPixqPP75uEvS|wlHRM>un9E;7w18k{i;TGv9H9jX-vF_>$NUiz>`};W)T= z7!;<%+l;5^v^sN4dYtp4#F<+h!OWjm?Wo}3lCV~R)m~=v9V^2Q_Mpo9^&?wi5Bo9$ z9KLLet3Iq?RwQFukb@@&Bd*vbe`hYo9+jpcVzej zYQS-=)@{bOBoi+^_2!TE0IiU>>^fChN(2Xt9FByke8Hg{2NZ`^?3?-Yu;`k^)w3Je-(S9@}z~{^| z{Ap|z+)l_dt_E|r4mgaqmxSYxfCeLXd#Dx?IQw449lAL!njRMkmVpaKC>5zjL>f%` zYQT51Wjl!k7($W2{flu!Nc(|Yq|t@r⪻AGP^c-?}YDcpk%{w^N?fp&2L$R{@q^P zkilBvee&Xm8E1w*Jxu`*?8%Wp4}JUjirvKU%j-C@Q~M#rGbJ)T3Y6c=iYmD~C3%i2vOX+a#_Zxz?p*Iz zWPDXBfH$RL;kP8?ooxMJJn6P=y5TB(=J1<3L)&oOOMO@jE)&R_PS)5EPfg4qbti`{ zqREIj7<*T)RNZLrR3(`#uVH88Wt&}&k^ZK{v75+E)Uk6B)R6d3el;% zDK!(1Q%Z8|Q$yZz$?zml9jP%l3+$EUeUfu#7eR3li?^EB9(stowMA9kaS!6UTeJav zj7qXiZ36?Gfn<%hWZHUaDMe`w$-0+jaTi@oh}vPIz$0}Dm9p{CnHLkr-{!BqY_7Q* zB02mtZk1;bf5Z>(d;R2f<8@B@O!8!Py$KJVpI-M57NhRll3G~D+0tA-br*!bji34u zR$z^Y1=aHqKjtxG8JPmYcfc;<;hQAc(xC~Ht=S|3xeIhb{zK6$K& z+B_QAKC>JwlSK%85>Re4nY=O5&$Y-!?P>&-0-C;E9DDh8IMJg~d^1(WB$H zk+u}9#u}FAp))`eix(@6u1YQc3wK#elH$rBd5Mw8{_d8gtvi24v67S86cyX67pB$k zx>YYpgcNA!C_mo)&WwX!NtK&O&6>Daz{8^2yFB>ZIAlfAlA~VwWm)*Z{(7Am?wpK_ zK$Rk~cUMy!n^1yVDRtZjie9$E$7;h;Wq0H56ng3Nx@4wWUTI+zW`7v%E;P$k$Xuv$ z5*ZC*G)hOQma(!(;FP2tB?dK>SkbeY_4W^J7O+&c!e*G+LF+yQtfG@quQs$FocAwA zJY+;%GDas@ug`x9wqJ(t>at{%SHY0*q;@~^qgjEgew-J~uI zzZ)Qdw=|Oo;%V<4UZqHv)=II;0EB56E6??O*HMT}42P6#a+_D)#@S&P-%3>3*Eo0D zA>5|NiVJZ)OpITGm>bk8jya#DFB$Pa!6cRC#mFJYTrMXv-!Rore4m6M8hKgQSQ7=Q zbE>kWKeqPxzy@3NjF(7{>O(M`tgB%P4X_oF@W#hldT8Or3JUujbSd&D6llCSRde98 zte>+yd_vCE7(ZRVOx?#JpC2S5%gS936KkF_o(_oLG;!n2K3t3~aqDkgmh0AJI<#wm zIkGcQm50ww;mN#gs1K23ihRZzcCR=_0Smp_9|avYkWB2OvFKX3NSG7$E1_`;$>9ZI z5VwWQ9$$?FpM_JOA;{omu9F-+tf_7doQ(4GDNRmqKhu_MW2j!$efr&(*E{ z`DrX~N)rjlF{@mU(^M-f6FuzT{|VWT+F1AP!u=Xpa)<;TwpSgAT!8{xM`AYF;ja~7 z2Q)I0=sXFH_O=voqwLg}xRuD>%&S5xjwAB7i!Yni=ZVXHvCp#!QN<#G?CE7( z#WQU)=nb{4z@SOuUh0AH<1dj+A8$EBHCy?Wei(z7f#~lCU)~AeQ`ZAVBy{_%J*r;3wIWxr>1zklN?-`9xw#qbDBYEgUxnGAHZvQ&=43jI?c4Py^UQV`_E2-GXP?qi_e(Ro_ z0C4c_X@Ud=Z$t9*zf11l1f%K}a?Sjam<^~{T9#`bH|o8y?qS@!@tnK z*}!}hrS^aIOxFCU|8NB)5(~Y-E`#=u%B#;l7ucQK#rT;W)&5be)FZ-t@}%fZZ8A8C zb&^%`-@?8icuNj?btRzkXu41lIMKoU?wdul_ke83p*s zSlHu8>%DMzVXZQ^b^|@yQK07tc`|WqYHX}x`D4>pM?&0}(Q+2_DjWqgamxF=% zyV|0cW;8NfyNF4_cw@_q0pdMdHd~x2oE|>tdzhnu2uQAjI5*Db=16hpn#N~ZXaN&^ILQUkzxQMuhZ%A|Pw^jYQez*WZx_rC=N|m>yC+wje}e?V z9d(PrX=ophpGb9aMxSLFv~Q%w%hjiZ^nJ5cG|}bKM_iD^>{BH7Vz%?;_&47=#x2ba znu)X=*IT2Y=OH*T076GfcLXA$v>`3t-k04`H}=&2s_QYI**9(i??zb}U#GK9SRjtI zG%ohcZ{q9hd2Y-jux`Ywe)r;#K)u2DpZFE5>3#*PcBqoz_J8=d6o~f4PUo@3q*$D* z>*)Z&Lr%w$yHCSk_xF=w+0M4r#AST&#rN=?kCMt=b_N8_qe#wxdsl!2#HJx-r-`91k4z0;jqD+b zs8(@KCPomC04m0Yg_OA$r9Y1lub(*x0=+voJBV_h*KMw4Ym0=G_3}t}u1t97pCn3L zFKZ)#YIxKO6aP<4^-%2XqF1xuJaOH&V5f@*`#hcSM=cN)yYcnamg(P}{1v-lTHN@sUgo=?K5d!? z>&LRP*jckzFnmk@PrI{Z0X}lCa(^YRD5G4h@sLfsHeNcLb3lX+kQ=hTso&DozDdT^ zVrX#7cZ2*eU%E^s(ipGWhfliWHwfFz$n8 z{q?^76pGnl$x#QK0srw&4-bQW9^WXIA&lrO5uGEiRw^(6J#S*loWgg(VgQ+s!6iptHv;Ib#VPqI`VM zt7<~vx2DNdiB&33&WZ}NvsyO$ANKe1NCldcs3T^R!R7LKgKK>o_{$7|FuwNbnxtb? z-~YTg>^RW6_Oe3Y9ObIsuU9IApCJL>3cl(6BR$6N$rGg&Bp4_?doFXiRI4^;vmHh0 zimq+GRCL@BKsm8kS@Xzm{q91`{$||2%?z6Ez1wr#u2oC-LEjzUn;^rl$*3Fp@~Th% zOEdRcE0?VkAAf~x!qt5{%3e#^RJUg@>xGJg3 zGcir)mRI5Jj`p{5S$Mn?Md=n*bT z1yglhgDI@0gm28~zuhpJj-Z)$9D3uM{S2ofm^U5;v5#)&3p5K;ubM`(}6F9^t@enuq~ZvMcm&PK=*xAZ)t)x^}6j>70B$#N}_! zaUr5yVka0-@TlcA!uQL@MG?7m_=t?mM=_7*=`hIDOM6tacde||`JzzFt-4h#ZgGHW z*vR70b)zx?kx=s3Zqz2(X+C%0lJgYQ#xdR@IHkno45&FY-vJrkQZLFms(gm+n)BHwaQFH4rnOb@xt&BM44%;iH<;-NUq9m4#O_%8I&-bWaOq z*AAD8WAMMC#d+`?cuz9g+IMWMia7kNjcO6!S*8uRJW?R>w%c&3%{V{*lx&peoqy$` zC@d*k;(+G8t*k86SIj`#n%d=I_i~`_jh)1=ClnF=&c{WMf%$fDI$>vH$kEo;pq zoW52gFUgVli||*XwEw7Cl)BE;WXi)XlGMEF^sE4*Gv|({24%W-(kgB`?r;;7j)e-? zJR%CekzKlD5+=Ha*jYjp58T)**VL8sE`YnU(}?OqL>;+2g{}Is>d5c3K3ImaQrr9v z;qzpM=$paaXP1J%*^{c73N+%X!EV7%R48LDd7`A~!1P-@rjO0uUEjJ5IoO2Mcu)^t8q_e+kB*Xu zGIljcY^ks}*k2w~%AIOno`|V+^lgB~M_bJrHQPzrbtK@h5GOylz?vND?f%fB+0Vjr zuE9+d3FO4{+{3fs)3s_P{@Okpg&s2+wzqXh2qHZphiYC<)+`I8I8T~O;g-)g6w@KB)y zW?5(8p8QCpz$(sJnW$0}I?G#qD%>0Z*MlrIwSikO((|`M@%}%(5%RS1gPi{P8B4Ga zb)x1WibeU8F7Bw%E7e9qZr#l;K*d-tkVL+4Bs=y`>#9UVxJ*@(bW{6y-r1R7&L>V2 z5zJ%2RGsqcz!E&3>S{meF-J&2>)h9`BR94uQ~p4Rn*ioYAcYDC@>=DFoz46ye$U+H zW>XCbJ$M!#bow56a58c&Hx#t%n(6=DJ_UX}L;|Swtmfw~*WA^PY*4PH_zuk<@MaAI zTWD@E=BX=@r8KFTJVn)1bYSt+`#`58FDmMJyo%NQ(lN99In3)F;TmQ}uAa8BKh#z? zPDe4)4p7}T{qW)MAKKp?ugV1uF$6@9B!ADt^3K;+vaPi;CrD==wL}~a69yKS{Y4Ax zA$;&SiR+^_49K4H2CChGJ5gq}-*dV7j40fE@&U?*ikYZ(9Sp`t{(orHQG(=gW#(p| zT?>qU)Vj`?Kk8dkO@WK*y|8FtPYgbaCKNc2rg#~MheatuA^p&Ngh+3Rg zUb_+Bv`WTtte5VSjpw4xGXU2+3+dq9QTg=TLrXqyj)p{rWwXSwyc|;bQLPs=SO;Bt zU4XHPjH;xhq})0N<&`^%M`h_6A%XV#jO)Odogo#(TMY{YVhTLCoX!psf;5ORUXqK& zTwU3s)XR0m%s9BN2Z!wK6d(b{vSv<<4={sCTI2dOS+SA`wg&Jcd&IZe0o5UD<(`Zm zyss+`0uBE;#ny|#`^+)Edl)kA-{f2$XX0!AWU~!KS~&DXl?4D@%je z7K`Cs8llN^P(J;`*n_Gtng`gfA-sa_W640zKOZLNUo+OPXA;WeoU|C}&<1ckQH^~@ zb`!KiAL73W+d;6(}3-*^!8$o3K@BDRqs zvchUw$?*i2naU4w9dc3kZ>kRP5O00j_r)i#?Y~P9e0Gx{jurU^>B}`{^|Vqku#$XA zsCqe(8((91(SF3-uQ4m*pL9Wob=vb);F5e5C4EktALLwFEK%Oz9)Syv73Y(gBT6~b z2jD8**xtFyb&NdX$N5{7m0Q92b6R_ZDCppS4~45I&esHWNPy;;M)_&6aVYL>HI5+8 zTE9Sn>&|3UKMWhyYFyK{1Bor-J)H@@EveuaJ39s{)+uAjXJ&Ng{b$ZUQ;MzQ_vi2` zFyxMo*qmOQU%V5&LJ*^lNDjBxA($C267Mn~0f#yW2#?uu4mCgK8$yW=;_cUFz64}Y zxpqp#vmsQ)`UdPh+S}1YeGhy9r%h1NOItnQfJBYSYGhZ7HPOSE&qWzdg#?U;*%-h) zU$MI@GL@Yx3rAyZ6fw?SNZ^CJdiz3q?-441|CuBOp#=60iV4dHAC#;KV{~HK_OwnLET_eOKMH9iD-83%GGHffVy{`LpH?nv@ zu3RrJ^7X1=j_#wlFifj)8T_{#u=Bs>i1>}hTL=}jo00=;l)@!+dQ_%@82^`n%+BLQ zm?Nb>0)3{3SfgowJ+r*+*)`<_L%`oY2L5p}n}0#~|F_hCn8Ep6nO@(qj<)Q1@6Xw} zCqDVCEK(ielArtOz*{c~w`eFmadV0+f;i_|Bz9?q(ta*8)I#MbQ8qH*=|gF2w=S8- zwFs3}9{w)3VR>ZaAO_D@iAhfyb-74VbywEV_g%!XUFqd|ZucTny(H;gh`|;p-`G;t z@Ar5X5nL~bjS5;1z~W6xe6J=AAB9t8;)X?`@pK?1-K(1gBD;-EycQmfdqga! zQ?wu4Xa|x7x8lApoMqxeFo}2+6MqJmCY~h`tD@q0vB42l9Ah^A0^yGyyY)>TdwV~W z@LUH5|D`;m7sou6B4($JO_lmX6FJrUgvEqw6?#kGQpGq+ zy~paxX0sq10xX7LUq*L>@ixyJBYQ;E*erwf$BARzB<0^=X#Uyz+r4j>UK-jAaG_07 zo4ck)PNrx$oci)NSYc2;?C!SacLHD-rwv-6?vIU@LDh{LFZ4s4ds+vrC8?Z;nk?Y%OS~%*y+j3f~fA7_04N z-IH(KW16pftLs9h=n@YXBUyf(NYZb1{egwsC*f+6G?tx$j5BE32c!E-vLBOp6*gUa zq%g*0Cn2R@hf6ju3&Vd*Pi-bfuth>fzh_495U&@zNy>Ra&ffZpe!uThT+>CDWoCB2 z&_OsgSY(Dl5iCC3pm8|c2;k4KWnX$))_E^qsppZnNVf)n`VRxDylTG_CAvPEW(u<` zjoj!v!hWhR9&N(+s^OFCH+!8DbK6sBTKmtv?Va5PHi-n;drH0QF3Ox3p%H30MH}6x z{X7~Fo$+Mt{2g;X-04b%1^<>XonM5)SUSR9*G!nc0pPLFEsDxb3F7HnS0z;L9j641 zXwy_XhS~e8?{1YQR<2L^dmdKWvZ(9nxk-}%pIIuVr#o$(5>LPP;0r|+(bek3K%jFjOC8bZL-GQ&Rb>!)hAQAEsFQUFB zLrE zfBCniS>0buL!UcvTKsuYwGo0LLTFx!gA{lUs z6KovWUeSxgxYy1bs{fAubIBiHR?x=)(1LF9!JHZ7K=MI>@LYKzY2yaP0!r}+4;bv_ z9aY^iTAXXW@bz5=)v~;Jsz;;66soU)CX83Gj^W4l!Oj?)_rHX16`7?unp{0lu+-$( z`FMoAk?PdH=A~%|Y&x~-XaAfxH6He*So&$arcK-MzJ&3E`W->}$I8A(k#YLNQLT{L z>bbtV!k}vSl%b5u!J=MTdd!)oq%6%~ipuzIZ95I=^XH8dnO(KY57<9Uu{g5P;Y+fn z%N(z!fCXC?x_as+tJM_RDtE?qSze#L_FUwNXSBI+t&v=d{EnsVT=K(A!XGuS%}W}R;?gI+AHK#y{zo!2sE?qw0xz!b6P-pJ|&xJ20Y3=Xa45gbH%hw z@U?IU%b1P|Oit6E5kBEuyvEPs=!142-?_0tpP6H?8XXi*uIrCd-`0W+SIViSukjL=yBcD?ei~8Opd8&X`=^Z%XfZl%xY1 zi<-C*w8j-tCq6G&v!c~>BlLB#laskWAuBya|D#&^wTaP?z3Z*oP#W<#_@n&8%pkoA zE!WqmR+3ml=*yI@S?2)KidZLMk*`k^o(BjjVtUYTjMz`G;GuVeKo(VP43VjbQIedXl$7|zk(=JzDuSQiS`y}Ce2dnIxTTRL*3?*@f zFK3|xMh!mOcZI)3TL@IWCjn$F@gs?Qd*$oC%%1z{j&GOuoIcK@|FYqwrR-b>lEqS6 z#@MKSsiK0+IksUjr2XEcKVRh@7eApsAVV;Zd+2=9`3ZgeG)y*(;(pAMwF9Bgg+lMLSdvw~J*h2%cBEhw zI`J~>nNJ0NRjm`ty_^BG2&=}cey|bqWUp;;051itBp_cvV0b(~9rA&`_s8djgR!jD zAWl1)(}n_8I|&N$s*M5f`T-1?fUJ>R;xIHuel?lkGXwGg{J81yU=^WJ}iu!k{GgS z5w|4dWheIT1g_3~n;UE>eU%#<^Vw4KYBNZh8>UTZ3JanmdciYf@xs;P*7#|qx~p@> zUBgI<7&%W^?>8C!zxCYTnlA3X(G;|~t*w+;_-kc|>X$g0HM@grRt}8JzLNiTB=OJj jL}AY7L4x6;ks_2&i*FmH#s~DFxMBcGp%=hlME?9=C7j<` diff --git a/images/no-cpu-performance.jpg b/images/no-cpu-performance.jpg index fe18a17475af3f49b20242ee4097b2a4cfdffe4c..fef1806e36e3b10e61332a8d26b324cdc0e28e6d 100644 GIT binary patch delta 11725 zcmb_?XH=9+v+j!|35w*LQ8Fkw2bCm}Gb56-0|J5s@dep}WS9|v(z{T=leS9Y2ZVYfE+`kF86Zi%DZdR z-_S%oN{ipn8!6XRs1#=tWn4i&5y<6@PGjS zC{I@ZpSey_dP_4ENthH^pm=;`vg<@21lb!q;&gD(c$uO?{oLhiI3zhd=!%w(xLfKp zm$#Bveg2+|GS9mk`&~X`y|C9k9JTkuieEo_V0L9T)n*mpJqb|&tT!dzOuHv7I6O%# zNDSt)Atuo>nvs-eA^} zF=piv&bPPfpRD*INV|?P`Bi!d!0|W9@yf|sn@E+EU=$ z3HfiqciGxvsBp`}wK$n+!+7}?3(AUld6jcq@Ks9=o9w@Z>~ug{sLw#opFT;t{bV)b zxObHJlCo>bfFs=x|1T+o3>fGjXA;%Mgl<#=up4!{FGRzMt~;aUFf2cyfPC?-aZQ|G z@FTo8a4LKu>h0w9l(;zK54*eBy^DisM2#=Cl>#o6HJ0E!c-HWjc%*FK>Z{9~^X$*3 zW8*E!3@U^O9#-A-_O`OaW9)Ro6qQ~t`-v97NBG6zMN|%OmHP}J|Ga~yVSVAq+!OQY zRb$o1h!>n)D|*dK%;hK`g(ekAzz1H=IgIEgLGuyd;)C&n420}ryy;=o6gYHG0|j8m z$s(L0FO}akZv^I`fKPp?NKxZTokmI7W+YG~v$K!>&KhieW{m=*gk=#FuTa2IeT_YM z@6_oqtMxM|Cga-}mD(Z^h9q|g=0B_U@2sX}(4ok_fZcl(~5!@THNA+OVvvEe+&`V%!5M~+X5Ve~DJ8ATp@yZ`KT zs68Z8bzJ(adajyGBI#)rW3jw1=GR)j5gGps6JEmxy1|XSnUP<X`f1N`j?9%T}3 z4r4){jh7!b8J-9kVR}v8%zMT{^&NvAtGyYKPJ6|lU|^o;@$n+o zjF*#d@QpC)%l5`Cr3qdN`OqmNa@v|6U|dRVgoiG=^SOZ%kk>XW6^ z8$4Mg?3MnUg=6Y5)#1?XAE_ung^vT%3GTx&EV(QvZTM z`Oz?UP&hEuPZlt1>yqdEf#TraSxRc|!>qvXw)8RMYR>B-9Xbq_&=yV(mu~YGv7H2_ zk3}mGh-7SarOLxt+8ZC$c(&jP4nz;Dwp2fKmO}3L0Qio^-yR!2=wUE zZOm#ys4tX&gJhO#3#P=K&Y^X3wTju&GjplOyvXC}v7>?reb4REIGhYvvcqrg4J`i(bD7WCe=xjtsStb+_wEMJD!aVm%_EAaC{QMFl!Q3M$w%*TLx(VAJf%RDGi zQ{^{9|9SO3s4q;ktUM|q->h)!+hw3YrWgPDZwcYAiN@w?xYipEJuQ6`ijZ&}E+zCO&%SEk2VU$~Z~>4qXz%AQA*6nQSH z+1s6O_=yP1ileJZ`>-jdIMrua62ffkOcaWZQvHSY zc$X0^1#%m61u{*H{}y=jBYZeHpA39a1P6Y={t&c!z3J?hFY_H*U}~tw#)*|`YiKBu zA<6eVy85nw=u7t^Zw^e-8zLX!)-@)^S;6i^BS zw;nbbj)6XT4smprgNC}-f5@jIyb=mF!DFf;(c3@B;pnW!yUko-6C-5rp^6x&=SVN? zhZEnYf;Zt?ququD%qZ+1wwr`8^{41?5$UsHp0RvQj_D2a%FLLx`FyhVL}Pv}WkqJN zBWP1WZx{uvK3R6Y!DC0lAQO!bikceV=FElPA-u*2>Hf#$Xh@|tUSEP*;yWz`E$07r z5VW(9e_mFOf(Xx@=lMW$o`MDS1U#9rx;Ge2AKhTO*lIa!t4V%voVgOST&!c>?53AoWp`?Yo8^Ppe z(QNT8Ia=Z1C5#JNxo>*UfE#T&4{uQ3G>WhJhqch*WyJ94=I-Hbn}gs0c75fWhyJR{<#=SBx`69U z>!KOBdS2wcA@!agzBuqyN9Zl~aZif+hv;=aE3N1_)5bKqSb*>}jgawvqKZ3I0wdcq zPS7(x#JWAvXXGxhTVT(Ebtg_g^4s|y|BoM=PYYH zc2cu;EPXumYHX(43F2G#v;+UZT7Noi<-4@pfHyfAM$nSyqln()Wy7p5CJSmVbgthG zbZxF&bE1@~f3Di5IOP0s+D*Y-B&YIzLnPxPeq!BA)1xzvW!@QF^}d6Ndhhnb9EAS{ z3NFcI1vllw2gwtyZ{{quQ*)ou*QZ_>xZGOF##Pztdjy|st*9VT7jIRdi%X`d91Lq3TD&|*kW$yX z^!0pDx$E$v=tb`lLk?f$_gAl5nNP?w6zqG5yPAH9{*e9r-bq>vH+sU4Gxbt+KiSfH zr1J9c{Qw=zL>7Vlq7&nS?e?Bi8?6GCOJj*3%Vz;)oDC*-F~I5u;FKNDD+jKhN+weF_9 zH@lC!p=~%5s!HH!_Lw=?5_1#CPZ&PoO{+!$siEJ-4(~NJAN>-}(yU4p^43vpKsIZ_ z{qPyV!v2}`5Dwt)QrJwqKrbm@UrP}^=li=~p!#>W$6eJVI;MLtmtmUmLVl_lRo^AY z9aYbHuVkHD`iI2vao+n+JK`Klu~mkidP?8LQ3`C7EZDrRAH0_R%(FLVWCK?ZQTSFY z0I*?%|8$ius4NzgD;qQaRYa4*h{+cxb}ux`QD+U2BkTIKOInQns(8>3AC1pKL@s5Ac0J{T)_ByJ3eJ7J= za7GqJaQm;5{pZ21eXz`QnB1b>y#5&M#^%pbC-Db!d^!m?6zdcIp$hI8Ym$MU(e!@q zv`jSbt_KRRi#8Tp3H{hIxSw0+)W##=gAc5oY7jV7d+#N$UK{+eYwjPPeA@T)(T5BU z&HfC#zB|g|EA)89>t(y3%@OUqV$Jo(81k~gi%;h&bZ=cgKjDyCsVxct@LGDZ@zNh( zrRyG-CEo&{}5NwOYvm}y}l2{KSiRCk18yk7y z<*kDi&$Xw$?iC1N(h6dmhef9->rD7blRap8AdefRU~#GQSCOLlpDHEy0~40F@97h4 ziqdUS;u-9l0xGwb^Lw=zGgNz;m+FHq5csG@d=H*I&k9-*v5qFpnwm#MKlH3 zkVG}U(YYU|E?m_hx&bpxFP6y&dCq_LOPmjM~__-)5>0w0rhN@5nwz zhwq^Z)Mr`Dx*=h2Qv21wV^D4n8S3QqJHPSW;{BNC^yL4uptQKGD(VX4XRWXDsm-3k zof@?yN*I+$&wDGc$J0W(l|uAly3q82D(&w%qq!}T*Hfza1s~7ZjHr`1!kpci>#!Ik z%{K1%i5C&Q82=z!V!xvQiIjdk9J^xLA*9_v(}UpBJrLWMyk0)_F)VVQ z5$CIvwxIsjsDD9|l1JYX_Lb);XMDsmSy=9%;EeFJCYIaJd;LV%?J7J1_e$tOs13az zc!%T*nQ*>Q`!dGU@g2I3ap%gqB5AKvFy+iU1YM5L)2hC!sSXq~4}&VNXy#`OW3&`T zqA8pm7ErXJsyXE+^k!Cr%+%kd0R_+^}8D02~z_Y7f5du!sdXL@=w5_x94fH~jWI?rBO4dKpp zSJ@|YUXwKAPNU{38p`%s#oTo7YRfXV`@uSrnW!E;zhq2R%tk)GQ73ZpZEEUiuky4O zl{1y}7huvHPRA91zWZ@>$rCL;PUgRUjRLT2cguYC#NAMUFMabNeyUErqqK~X4vxj0 z=Q)>|mY*5p$!=)%%F6<$3~OTjD5MQ9*k_YcJHPL+HPG;T+`AYd_p@x#y+= z)y5Ie2&Z2s3*DI)jn&CFl^Uf+9@16SQ<_&+3-TY`ksDO=h5}?-3IKU9+??*Cx`>it zhCThK^iTa7W*nF@;AxZVUD>smbgu*CjkV@S$k6SHv0ie9V&_Qlns)iTc(OlTt`t}5 z+wV$~!WmlkN@hp1BV*nnn;)l#aP_zToB1FC3u$+5AO{cnmI?-BB)`bC*7T1)==~js z@gjOVZA136vjK0>rKwKy5y_6k>!#a8(cj4K?KcxJ#H+3* zA}D<|hgsG)<6P}^|Gs#)Nv6xBbU{|JsKTLBN%rh!>q&bvv}`nxsrm-L3kC2bFrNgn zgMM=-CNY4uXONK^k&0tRMgMzZg>|||=bMJ-T0Jst5iYNs!oBO0^l7wytv804=x;e$ z8@*`kEnv-GQ|8L_dn?H&VURKLuAU-UxbAbBS+kv}TxM}tIZVacOkSsFQwb~LD~3jE zdGCa1^|1WL4NbTd|Ovi?2X704Nr35x40q1$% zJ?myCsf&ygHM?)1l{3$)DkJ+o<1tb2sIb9Q#L>c*1alOGyV5y=`c7pxrvd6Z12cP3 z;23is63ZN1w5Ti)>AU6EX=q|y4ZDqHOixWh@^&ErT@BOsby?DkqRvC#GjVXx2v^_c z&tq0_t+@wfneHRgi9VJpizj1rC(J}zlp-F-xh!}FJ7)ztM(&aGvOWG59W-UuNmfBT zMz5jKT2&=^AiaAl1My3-3Lqe}ER0Ym2T9}#HSQYsSj&FlUHdk`F&>lYhESgv79`#j zCH4QU=Xz6Pb>cIdyOo%lR_g%CNC7Q`w&C{S#D^q%=2G);Xi2se!x!) z|8eIT$mQ@}%{Hfx?L)s}x53S*D88%!mnyezc}3t9o;i(w?V06ls;gH%+~xoh$arOK zm*R8Nfc8&JwDKR0FTOiI8hE(hp|pWCo&5GMP)V+&|HF}hrBwCH;?5+HR|wGyAuKL-8nij~DHM z^6J{R8<+~hQSuoIi0TS%-epU@R&mx7Ni3ft#Tspx`xC3ko5RZYO4)&e__CGG2faqDvRD2la7=O|Xd( zi)z-r8jV%R9l;Ww*6^a5CgnM<2Dyjc$|W@uB$8QR@l*Gm}jVEHt!&Y#l z_l@RUAl#`t1bM{+y@(Nbv0d z(ki*Tyc@9|S8;w=3hRX?w=Pcy2N__K(u*42BS&LO2Qp@Hq+x9JUSj zRw^njYT_bS!S=R7%K{Y?fblA|7}wh+Qm{-Qj;oG7LYP+{s=V5Uv*YOeEa0Y#QJ+?+ z{AQv_0qRViRjEJBq@1q#C8Ly_S6s45=mJBBC-{&EHn13$aRtc%4*y} zR^fllR!%jZ>Ve<>QL(gv{D_jQ?xW z$~=6E$C>07cbU%iAYV1Hr3KYOUsT*kX|Vm#pu$!j^;*fo=-VO&Ov4P=1jwJkoAE8t zi~tdcK^;R@@gN^Zg&n}}tso~)Vz_$czqY)6UG*Fd-yP^M3Ey{b$D|e8c)A zO&98$c0((%Zlbc~e>V-usYomD0vOXDzd{F|EM1NE&EGG58XWVQ1piA;V=QeZ5GTt=2}ERYMpATd

U4J_StHH7$A`C6&tvMF zzRu0&a2T4iq5u~O%!8hbB7FU?H6xUsZBWytTkqTUw{7%|c!V4zV&h!7;3klVCS!UU zwp&9&mX`^XX&i%L9C5K;R2+=Kh1KKxRno&-nrbAZfj?4sb2EOC*ST$V&RI z^_^#gIft%g5?D+R=3Y`B;Wt#AvV9~YAecHiZ&FvwK^je`eeo4=U5F-gTd-mnTp(3F zLWEY#{&b{2BeWlDT_t3K?zvXWgY)Z$=a1W;GG5Y}gT}CXEUoieToCS;5~`d6 zvbU%0sS_on{_xRIwJYqc64|!voBoQNSqFk-eTC!Tp($o|8)^Fb=1u8v`*bdrbPhc& zzce!pqcO+pmn2?lxsb}{lU~tz6mZeIii}}AABS!1daGy1Hauryym|oUA^eM26^fuI zHaf*wSx8kxxC1wL|D^3rsi86Frd4#d`I#mexUvw=IQ70_^|d1&%=oJ9?rsXvLGZYnS3m&OCd{KI&Ux5ey`kD%0v?)MtnJ4j5-7a4(sHS72KZPii~z-~BcF2VE-&z&F>+T8^m3%<%A=Dc8E8cF>ui>i z#yvb#iN52sKjmF;lcj>9hoq~3C*ZaqUFyCgsz1io4C`H>ar*6Ub@LtX6o!oepu z(D4E`86$loc9;SyACV$GSYV(yc0b1ENUGIqJ$@raGZC|)TAh*|cnkC_-q$exQCO6m zNk>P=a3-L(yDK!Q!+ztmg7l*)StCLLbQymi6nS%r+aebH4j4L>&gZ&r0b$c#Mx^-P zKO3KF^~tUy?i}q}fzVlS7*h{Kd`AH*U;jP&YYk}PmYNFEIL#YxR{ z(46maW`Xq1YvM6$lLz1Treg>!2BELOVFF!5I$B3e9$PTUe)KQOUlG_6sXQnk@sOaa z41QbTf&%1tHIC8!zzB5n;tje}ET^XW=PxN62YQ7VfS4n_f+1bIZ4gr*1+ z{`1C#Q=D<|;2R3KtRG8{d#JIFcyN`0$hQ5*jBFM9c1{z?KqGE+<=runOy;VKbpDAe z1jd~_<`cQq z1DDQVuT&2W@F+XkQnQPOe8P9t^r$F-(?=q8=#u$H8Dv3#{p$wWlazY3wGz^85Gi7| zxAD4f-0cZOm#J|PI9X6PW5)MwuBk=6kz0@fAyumAOt~N`i#A+o@tl#}ZBTYGVEWk= zL$d=e3b4+uu32Wi_#mM)q3HAF%k++ldg8cxP} zfzf>KqJ|tfQf0h60(>C;h&AC~BN|hE1!!E^FP-Du14j4IHj}~!ind%LRlGi8h1rU$KD=FO3-xw_z2O(gaRC? zz`)k_TxL0}&n*|74sr>>%!xleI)W5q^}a@S)D_S6P(>-4JS*+i_LBW9ky>FAH_I>z z_9IW=At<0-i|==(*b?Jzd(P>Jg6z$+_6?+o#|;6wY{|6JadCD?iSL^~p>!W*6Y|`3 z!EQ43)@74{W})N35>y!7O4P;{}mTbz032wZpVoM2~c z8f|>W=Y8-)Ox7=vF-vd!^c!Dg=q3+BP~tT@7-st3mWf3?&E)KH3!`I*nHqJX{j}=D z%4C}a!j*)bH_Urfue6e(Z_U3b;DozjnH`3(XvzX-g0Ej;b|$)li|D{O0^OZm_h+IT zs5`wnc$L^vj=`WS%6HX@mfNDeGA1N@&3lFDL~NT2(f@_>WT^qY;+ikAm-SQWaTReV zvTcZfcn?Is>0USl`8ja`aku8rV4j$8e=By_v;)ozG@~zxqF#=3SETJF80E&d5Fc<` zQ-}hF^L){reSP%)5v~O$4j-8kmV1^C_Yf+C9qSsEGpWK-H73nFVps7HB&Lb=Hzj6^;zmQIe+sZ*TSE%P48@s$=dq9{n|JG&$0ftGAzw= z&O}2|C(3@@M;sJH$aagzIrncg*SHT|lkuZ~L%tbw{6*+MUQgCywo_tw+b# zeaS7EJT}?(a4VD)FrWJJ^eEwFQ?`ErloW1$HXgwH^E6P4<&gT)5AgW(ea4KS zuH){)I`4S>W9*H)SSc)n!aaDi?jd+LyS4GrNVQ~+%#-GD!>~baQ_>x$KY7l zK2)%Mof|^?duDn&Hk2()c62x`l$&@n-&IWdz3jzjZ#^BjnFJ;1u=W?N@)xy|Ts%c^ z%dn1{;u)1eacpY-7UN1~8{;61(R(D*`hJ$-olLV&Vk~}Q(hL(5w;wni^PZ8NXS5-$ zT-3iu{a_}1>&eF1Rcg8In0ELE*0Xran`P0}Lu;56`!YF58u`l2n36^&{=GeU)?l)| za3#~-?=W&JP!0#sRK%ronJ(ta=g`y4#{n<_e8BiWZ&ubmSX(`mP?*t>-LH~-TpHBw zCCF#!KfV&-ikt0KoS%_!*NI%;Gvh0FL9Y9~K%FQ>pMY*5jBbZ%>xjIf0oxqojBT8C z>H}d)Y7u{rsNAloHC=Nh;0b%lz|@59JoPZh$2Yz2OQaUNmgw{HJh~$Gj7QP!z_=37 zQXuRWfIlHLI8J^idE~`pL$#!yJoq`yV8Ju6?3ttIP3lS`?d& zL~s(z&t;bbGw85mw!(9C+p*7QwW*;je2rwLp3U%mE1wcR4i5Nb^-OqPN9x-j;JOQm zSI1Lok?;=rZt9&dcvq9mf3fbgER?|`CK)`LxUZ~zJ6#1|Lxn4*HtXlU5?Y675I-Xv zaA$pNZ5<$EzO4(xptB+oBj0j2DSR`}{Uu=(Ct@`=nb0g)1t4j=tOeYD;{&aS z1^==I+9nwPyN@ihR>8H$5RA6Y*kLD+zE~fhk_I~|5?9cD@OFJJKE^psJE`+2RR81rd{#p&x>uvTT# z#n#W$4(OO=I7_%?I61ML9@_8k`{BF8HMyig@_zo1OqnMVMm6#|lp>+-aQM3i$~t4T zcnIB)&3y*tqAVEI%h7xaWq+U!(;wb4r0)318FTwtYsgNXNir02^IXnKgMBalkWlB%2t-o_dF5+ zl|!ZtB$bYqbdyseFEqrYX!;YC#`i0mX-F+BYJa75R4P7y_-&d6%0@>h#+o8|x}FH; z>{#mPszTPwDYTRxjP0`oTm*P5al;txu3Rd_Ho}H+wVVrno2WUHQL3=c&fI3xHwLqG z)B0~J^*4XZgYp-Z6oq$G_OUU~iAKbO`dkaqShmx(1aq?5!;QF0Jr>VHO~yVMcae%o zPHU}^Q+2aGG(}>VbS7sjWNms}TLBXCK=Aw14}{h~a^Poagd0a6;nK-tQwxd@BPLY8 zr7TmEu_f4Gkf7qLm7k5Dy*`#bufEk(w&UHxro?OL3Zqd$$8ADfv7WpByd1nKTK9 z>h&apm=(o##t$(@Og$5o#++#;_D`&G$$Gxb#WHzQQ{BlWx_PdbY^oEgJj&-zty?co zZ5H}#$eLkwk*LVUI5OyrZCrR5tfq-sHO@?VV`|&1Y(bl$%lVk+Zs&YHEz`tC$Q!H3 zmJ%$ES5u*0A)Gu$=F#+i`|1yF?e5IZKP(&V7fGn<4OiXWkt&^$XZxlTWr$7a@hJyB z)@)q%ewm)DcVc20QaUmhHuHrz%RW#rMi;B}CoYqLq)YBB1o5mcrR!y!HfSUSalO|V zoktyOk+Welht6U1^1nkMt z3)g&km)J4q>`Ph_M7l3NswFi>H23J@o) z6uF3gB%QFWfP5*M;4XT4z6SC$lN6g**0O$th}qDRE!x6dQjYmQopJ6pDp@RxF*({7 z3B}IdJzn#Osfdv4d4A3A4QxXZM(M9QTzKBsDl<982|PVof4IF%FF;a}`JyMQs&nr`%db zB|8so!evNa^(cJjUdi5T>BO>fsXqJhE5?hV4hiAg<7vq2pqRe9L$TYim~GC2xt znWy=$sgKBz%;WCbKeVTe)y;yXf+?OwF5A4h?{y{Dy&@cEo$W|!kE9*R88!a(F4)|w zgs`mAiREcVA4aHk?M*LOi*2_nF&Oz&BtNa>e3n-IHx``2&0IU_JDT} zO6vXneej{-v~hh?Qz5ZHdugC@Yc$PQ{9V1WA1d`3A$%Pxr&;qN|Cqv4gI0 zT1?zsZ4R8R^P}jAq7mG?ibKW^7|a;WsNgB*`^7Y~SrPV~vH2xc$KuujZXdLXXi+N4 zPa-$;7Ij=2soBZo7f7+#i1QRA7Y^TQ@03P1N8*#R%WRcWnQM#Xyx71gatckt~92VgpJr(@)&`-;JlT(>~il>d8&;`qM zcpW|`OZJ`QSiAdEG43SSYX`8~u@wiu!-7A+#H-}U0S%03!|m~CKtFsA%t!r}Y{ul6 zO%d^EcYdi_0rGZ0enl;8$pPdV{}Qc{qoa{+O=^!Zn`v4DiH{?{ahFC97(u zsS46Cc!dIUa2WIKWw(4eP3Q9beKg?rP^&&_51vLq4#E~$Fy~9AO+gXxfoHR}gQ|O5C^T?( zt?aV626ZD24cs95h};{&oKoKv4YX~{W+6oeYnngelP{c~9LR#e#@#dRJQ6nYFA-(N z<1_cNe2f>-0QTY$KVx_Pnd_6f_)gyA8#}zFmE}B&k5aBfwPqhjD^#-O2&d(H%!8`m z5H%JyW*!-zehB|4NWVcV@EU0Jo@iPh$EC?zZq(o{4x-?n?m z4LQsaX&Jd+md<5{g4=Tx?I?+hu|hLz-JqRABW(DJ?HxX!v=B64@LuT<+EBbH8#`EP z8Oo6l{mW7Q{^X+bs`07;%MAPJzpVy;3ddhn!c8&7Lvqw0f-c6C?E-lT-x8;Vf-7&# zZ`uV7^A;O7#p_VgIw$`qb{yLfHCOEaP(f-cSm__MKo1ftc>57=jBOB1>+LvLDy7D~ zhK9HlNIfV>!3R4S!OI)QXkh)>D6YSG43g&SdHgocL_l6VF7qqNr_afTne~bgklXW# zQxRuej;n%T`$$Geg*}4=e}e%4S%=Tzn}YWo5!?in+BJ3M6_zv__wwFF*(bb3krGj% z0#2CDNK=*#?mb2DMO|EWt5S}_8c>iVj9ub{p@9cpB*+_T`x*cKuI3-ci?^#*dWi-W zXy*5#ts9Q0zM+A}4`_fEekse{&HesfewbZV3~^&@kP&#cKuSVCI>njASA22VI>Va| zf=H~W^8M1;jW^gD2HBIs1lCzq(cfpWdLjq;qSkt}8qC>I@+gX}8(r=H$E*;Bc+dV}KCZq$t()nl}El z$#HL9e;vFCFwnE5be=SNG=?&>*A`s6KXWZMm_CReE3zsbcVkW%QCekZ_b681q@kj$ zDpoPcG^!_7n60O6w#l3*eEdPSE#*a4>~n*UcAhs5W}}Pth(!pGwI)vKBZ*v3>u;3T zf2;YPU_CYU^J&o4l7|SV9UC*taZLqbv>}Px;@z3E6l|CPLdU;7ee~^fxk$s-fQYD+n~?;bXxVL z$va+(yv)ImRaE8-ZOZ-cT!kes*k3k!&nyt?Dlg8G7!MS@n!T@L_N4M#i9c_1(PyYj zHESRf2uqfcBur`u)MUKY9_(0_vdw>d!B7;lcFy#=lN5DGSJnUJ)U{1KaS!6ExnGpz zVCKU^Dtdpvr&Drl=$v?PZ$W_RKAx2L!m_BauNs+6jIHI5WQT4c!`fzCfy}~)@1F(B z1ZQfk@*yMsdKJ&{>njHwkDy~?ynK2;6*CNO|R(z@?;?oF|RXPrtS6O zWiJvgyOBC)zT^^BYKPl8nlH)YaCc&bUMgS%v?;ABce)?f*HlDqysijZT*8@q6ludk zR7cUZ%x2U16Cj%y3o7JDQQs?QV!+xUY%K#DJ0A8A;Wpd|Jwg}^K91@ogK zr{;%OY{q7z$JnlCQ-&#_} zm+9!>&_8FY<}0@FQj{Y8*oW#T8X(l26vyG>NJl_9;Pd!Yzf($SBIz+i*?RuaG56~H z7D&_r`)Vig@OS)r->?NuxlnR45HS|+O$Uc%fBwRxJx%U^|5sl@zw?5B{FoETN%u-a z&Q=l5%uIbvhaLU^j|;BEPXt%x$NQ&%a^&g|eX44@nT@C%L)-B&G9~;6?2}`Z#5sN4 zF;Is9!iZr8yXT=mE2?D#{AD)=+a1kD1JoIpBpPTStym8YH0Y!tP0N3!5~^)UDWHPt z%urVb;C?9V+cDDx?QzuQu?;4FaXbF!OMw3$*7$!}0{lPu6unv)M?+an&Db&BZi%8r zDCgL9^;H@n@5&D;anto;g?Sv;^1Y`L%}Qk3_o1MF!8Qr%A>$o1@GTZxzg=hDa%fAX zay6=m26$4g5=zTp=Y$?#P|ZqcATu63j^lN_tT?l|px_IcP7meQl>H5@{vmjdy&VYn zmMNQa6ub|5$46BR4M(AD8o+q!3$x!tJ1J|A^1t|oA%WO*OX=0P(^`5 zR$2MM*ZRsjZ}{2_UW<>koW_NY6C6h6269{r5N@ofjUVI4O3HG0&l4nUDp31Phi&JA zwr8(qR%$n7>V(>)QapAylIBScY|vpph%0kHvgKl|EvfgN8s^4 z%ryq7!qnFu#J+}4?jj0Zwt-E4oJuUepg5^Aac257&?|YM=HSle6ZurJhekUDm(DQe z?VO)PHT;YrRFI3H{h%P*dDP4l#+976K|i&Hx=i=Ahg{v&A9c`z8-)gyJle~S!*)gk zZn7xZ??>>h+l*jYLBabsUoZ|S@IL|ipFkaRi&S-q;tplwR8#{q3NEz01+0gbckP`K zPxG-p@QG2ky>Tmu!n++avYlQ&Ecw7PUfi?7&z3TPBTkvc1&>BZP@7#kn0l@7>|&kg zsP4Xv8{W~xxJH=UJ|r0uvSl*DawhXg{_jL?o@HLETY8w8J4`bT1uH8Zn2c zZOC%=GuNY3;^fevogD%?eLhF?TP@Vxg4-zA;~&94;!h3a;*JqXJ+a2ml$@MwWt3>5 zEA4lfIj-GjmlAhd@)O5z-OEP$V*=VY3}!{e_3D-tK9SHS@JSm9*Ts62fKlML@)PMH- zkHDqvSZ}jUHfGfRi#+9LC_>wKZDffd_Kw*x6oQPf&|R!ZKz0SK=Cux>ffd0er^_4c zs93`>L~n0hQrSl;tsmI8!fQ(^wKX=B7IdfrcOBda=5NjFKWlMyyolY3r@_iF9h;%Q z)wFW!OeZsm;_uc)4*tvBs4pVw>lnJ9*@(o)W>CypdB5Xm$2EW7^|QFAxc`Z^Ggatr z=A`reFIm$Ij&`!+cbQI=^YTz!^N+u@CPa)zAQJDpG-02PE8AAwEm{fY>a8r2Uvl5Z zVp}sg67!`QM-X%;u}+<=5xn)vyed3N)n{`v*HpwwPxliIJEWZ9Ud>%aIumgJyO=VMzu@NlQ-0MpMWw{7dl#J z>9uG;{;9h~9=w{75qncpr5fIZO9l#fV**E3ixop)(x!1=c^ZEtv*tqqtA!N{8LTR1 zZT=g-623;biO<+1)0ZIP#86!kqq9xtuN{G#5U2i4O^>-gm50iRPzkJDsBTQBR7et= zel*Tvw*WORQKZXLLNn;y7M%p2`V6)yk4M%A)QTd&w08^J>ecl3BPqf5AuK=fyx&eL z#8s5i8w*1o@v>hXDWMR&BgxvCLsoQbmdHTfT3_uljb*x<{43n-g>owm5@3w6c`SV6 zHZwQ^cl<}t{T+mf@BVm0S?ZH3JM}znvHOcsrb0DppkijHD>BdcZ5$0X+Zs*#lVPOb zVz2U(F@u9$wY}MJV=L=&8_F6+d7ey4)6@kANn&EpU))YcqmK7F3WoXy`5cS+AW!FzI)*8o@XiA9M7wQNv%70ZQBh>J z`#_RBl0fLYr!p#zqeDXE@?iE#6cZ)jejL>BXb!ypPC$(A@T-IG+o#N1s9mo6p%Aut zv&f;T_N(<0Uq7;|>jniGg-S5F3se>?khwtR-x8+Xv)03DCGyv#(=~)QO!4)l+7=Yn zh)g`5-{39wRv@f!HLBiF{`?h>p<&RWUxGq6G7&O_fd0z)8+QTzRg_(A4oRnj_n?<_Tt*Py2 z=p=1d-}u(>jGCMKx5k*~QaQ7TPw_FtvhQ?7e?lRVOJ6@{Pg>6jcv*alR(7`Rck<#r zs^A;UVE;t?PEkqi^T?i})Sf&CRxbP%eK0JwHu{GRH~;)}AR6%4gK2WT$#*-PF?=}D z@M&HSlS^c|F$wj-^ziOLj-wi(EH&HvH<)rIP&$-rNJ*b`7QTL)1$^N!_$U4Kn*#e& zApK{x3euk)OzmTL)(+N95xSq#ZgpWP(pMw9%h>3P*|WXJ_}}W2o;2HR(-a}3__qmO z-(D%L#OtgVTqD(XWsP~;?;C(XI+{hbL^(^|FEyO*Dt%l!L+1{8L!F6j4|kuT4?kKK zD}aH3PtZ^G>gu5_#<3whxuRLw)z!M^54pB1S*cN{42EwTv6vaIT$8y*b_#C~Z8=Df zKPm3aRY~MUJ_*7wbZ|rbdd96kA_UnunxtEOrc&2#F z#~13?orsSlT=W(o=KwCC%>>uLqy68N z2LH^uptkANT83m&hIWXL?>@>-=RH8G3hZjzL+*!h6W+`00Lfx5OiPkW!+pW=InBZQ z2EXB~Vv&UhZ)UuNrz3;>BMk~dREsm#A?vqVe#ws_^=t`-N`(xF+dX&5`5JzeFDHd} z_dma}X+&ZAZ&>da*6-L1{u31wYiW&d)2iJhsN-d4q|@7F<7+&RH08s?_Re&H{4BCl zRVY)Vo?@Rcy)>Wl-ae_1yaf$A^%29af3qLuswY0~kKk3-9Lp}965zS#~GB$I0IYZs!tQqKt4HS=(h#_*Y(vu zl!z;S-R=Zt_KIefkCl46av4@++X#^XUJzf``EvU41j7;H?`#z{ z7TbGw=YOuFek6?{c*gG(KBd$sb9ctQTzw6@Y=>XKwMH~m9vBxj_-3iS#N4c`EGBOX z!C#W%Ju66ZQmP(h>@3`kpWYfepJpTxHu5`Hq6?mtjPhno)j^O$!WJb6w(qvw*Lf4h zSs5W9n?t9Njs5mG?>X7ftI^O+O(W?60pId=QF$K!g3d!1(VpAciU?T!-5jz~JIP=sNGgJ#Qwq1MWO*#wLkxtVDmP?|3irECG=Zxv zy8qhqCss9wOYSejDkLsb5(}wuC{uevq$tNAXz7;m$tfgncUDC&*$PdVUt+WEo`V zpUZd764DxcMLAEc*`q+c17-oQp2E}EsP0jysSmxH)xRaoD?PM-bCq(9BAWA<&i}`= zbU$vr^%Db&hRH3S?QP~xCqBG*;m(bhINTy|eyW#XdcqP1aL9yxHA)|U0qM-C?an=m zgt5RH*m-P{1Ceh7*x+2E*OPr)pJM>x6n|{anpI1K1)saoV3)8<9rSN3|A5dKt zM+G3)=0rr?pDphO?V~22WU%jTJ~47qwbT}2m1@Ty%4hw4blnekFZXAZ%hiXey3c8H z%*=idI;8Nj@|gcZHvjg_+x2mx+d341ow8l;U%9JtX~tAz<1)W@6ZVy!M_%TF#sm{&x4+or9rPVb^VG4(~2s-l5Se&~gKn!+)xka4gs z9fL)B#JoyIb?r7SOw-kO#0y?jA%kpSVCi<$d^n~NW2Dz`*sJ&Ssd`S@Q+KRuQFGGW znDEDZ`gq>D5`Lwf0olA8$n;fhT>i{$S44(RmggKz&Nh zObb~0=k(|3tMb*By4tzkE-R#s)8&{Tgs{_4i*`1#C-e=Yf%39XMZ9K1q)T@_U4@iS zQoqD2B|m5*7WMW8DGD=ef*+T9B%4ScJ~`~?Q1PFgW+#CAi2~+(xa@_-Iy!wr#KgoT zZ{GjtwlUnf(&8i4m^I$Lq7~@NE`ED(<|`z%>xTv=ZD3_c8E_496Ae@hfF}lX>7^%_ z4p>Um9vWc1Dga}AVQViKlv_p5UvZS~y2F~`mkO;ZM4_0m|t4(@- z_tN+D=|)D%!Jl)zw)yDd#b9U&f)p~g{3RPIJ;|hH(HH$*KT~=G-Ct>)C_Wzod-^PC z5NA~Ov{G5s$sKtLzn0S4LvJ@!dIsI1!4h;PV&~#zF)3!DmAZE-yU>5x;-5pN;(>ufzE0 zZ$pswz=VJ`lO*3Q;(H7-m>!VaqdB$_9QQLP4W~=+XUbSLt^A8T{SbT{C3ZhfFHE8O zd-Y-)k`?2q^TfdVs=~tFc87SJ{5UJGg7L!&j+ek zo^pm4O)^yNZlMv+OrmLHeR2b-ifUN>=A}J!U+id~FeXP2DOivm>2DH3az%p)RPSOz z@PhKV;qsUY4bT^EaLZAaJxTuKxG&H^vl8rz;LZHi1qKC0YD@HbX~l1zYNFF z!0w1K(xrm*Q5=ZXwz2U>DDkaoM%*c}H5%Z2i|MRkdXz`Jo8ajNKT1pi3B~Z70~Lj6 zU_PheY}H9afqRF35T%9tkL7aL#ZR&wGW42qti^lAzfZS)BrR}`KDZ9~M7o^Y`HMT( z5?`nl5-ldx%MybfAhbar9@x(g?)|lI@c-?;QD;i*m=1uGAtp^|fcsr@nNUE988VU} zX8dw3>L9T!;pFR7Z`pe4CTl0}nv&_ka<^hHjeHxb$(#WM!wxXn=HEWp(9}9y+B#0M zz{+Rq-hm-rV58Wn;^tcoYb46mtY| zzgTzkycT!N^|UkL8KSFw+t}_iQluS!6I~HpT=zWA-#4 zEbpN(uc)mIm5aL4toYO#a&^-RaBq{)RnRarG@P8ysw=qFgl%?%UXkLd7o;)Q`Zqwt zk}xWvx{&H2k=$pI=Uu69Ucfqlk6_JHOCNXsIw`N%TvRrCOr4EjWIALEbt4Si$1{7| z#G)$omPyNKnZu2%k%A5TRuf2}x~wqFx6N8sHC%Y7{2hc`NVXsZXvD{BT#l4b#{u7( z$E)G4J8)A6&vm#*(MCuzTIdPS#OF$1k>*hnkIksm*P`~;;>E5#`Kt2fI5(+4cxZi_6 zq-eDiOIKopk`|?gzl~`+Vy5-^23bgSW@?j}!W z4WdcdkblwDjpN-ik_){i+5vsZJ+x7yH(9HDOE{B=p%dg4r*0yR1#SoLzXX1MBZmYo znp0L(_a!@Qv&9p<6H>CElOqhw&=jPx#A`HTX7rHS%*s1)zn|dEC>K|7e_*BJBcTZu zdv4>d&L3q1R@{DE^oMmS2T3$gJ+jr@%yes|r^*ypD^je3@$*{MmE|2DW8Kk5I&E{~ zPs{Y7pb08I5L;g_N=_@LX%iD+YPko$2G9j+Cv&XgZ)`9;H2JV{D|z97NcW@SuvLTd z`m0MNezNP1v=l;o;wImdg{-fw{$}|t&feK%%YWEd`EuLbN>%OC!)v)^E|c~|b%&9z z25|2-c{KoZz3;Ye~la=Fs??eyyVlz9-6SjtG_?Uap9H zt3+y$o)I$;jdO#-cW#bO`Bg=Ekf=$xgW|G!Ud8~VG4;l8aLXM&TGq{=Wth;fiU1Wx zhRh+^r=X^sR;)r`NC7Ub66O?_sBn7Ls5!rOiS#_|G{% z0f(b*y7`-p7jk*5dUKspbD8qbZSYGR@&_{#(Qvjwbro zQZy>*HK%P1+~VN+1``HN*5|X>>UnLGn{in516Z%#2ew67abVl1s0ijYqRXXAL^|p60IsTEJ*y48X%keX+mu&Q&-{Tft&+RbhBlg zddVZa>mBrj$vAIZpOJYmX)==Q8a>|Hw4I+~#bWh@+4p}^c(UxSD^ZQ%*tOi23Mu<8 zmioK~Dk!`9#Y1Lmv-FV7dj9?gE_cTh%|na49Tu;$;!l(WW5Gu{Lse6-DE!Q=Be2@D zX<^KKN;=yu>ulf)8AV$m4sW^;3s2i1nGN%TNfZ74eD=NRE0`f|Lb}1caCAwiXi$nX zIa4iN#f@jtmkQkxpW7uZ{lvq3o;PpUd8Ae=F|N&T3v90_QOeAa$Vyuve<<7dIQZB* z@_yB4AzrV2)dhanL58qNk2XGEzlaZ9fDtQ`LJf8jNykJ#H_Z)2apbf;py9%dn&I!< z+ZN*Ge^@ch3>m`1R&K~xFylRz#7sX}!NQI{Z|x-aT?q2Hmpp4dz1a~Xp4$1+J>r*W z=aYTSXJ-`(-_-dXlkd|#6e$1j)sVWyp{Y-vJHm;?ZaWK{nb4aso|>-z;_7reYj#HK zL0vjG?eurk4Ut=?uHhL(ss6@=Bn!A?8&N5>jw}1r209ukWWXubf{lwr1Nj)&oUW`Pyl(90MP3}I?0}D5 z=!iEg_fA?782D|Zm~|9-XnHAi@$lv>ZJLt*3;FxU>XSGXl3Hd^-|HSMvTHqtJsL`; z-Gm`qxGpb)H)SZ-_IV@YP0B`#p4u8)>8}P&-GDYH1~hV*l~G70CtZw&HQZd+(J8MW zoehfIvYe5L;C>>Qn3sAx2v6)<)pUgX()t zZ=GY3*|77;)-huyIlHQ1&xgJS44K!3{CG{!Pmt(0m~;ILdnyqAlfmwN&0cb;RKkp$ z<98!lstSUGoF1XH~IOQG)!sMqMH~$In7*zHW>e{HWyKR7O;+-;dk&bDxemAG;2P$^dj-1Px-GX1)tUb6BOAx)@5hhL&cyKD%v zjKVZ&Ejni^re^o$qf~qyU(Uhe-T0xYBgrB?dGCs+ABxo@<;UAdX4$he_xeM!_lYy| z;tt}whFr0gRuoIHCl6le-+c<<_cQQMn~kd7V8&r~Mp^{|M@a3-Z-)FU+H zZJz_R8`+C}c8T?#O%NGw3wrwcX5HIIHtxu`k>Exqk=%nkR3AswiP33Pen&*98vZEP zqnX|&j$jT~#5;FUnxSHlSVJ6bt13rmX$}`rP%aIQj(nBzct#r0?;WJIQ?)EwWCB*~ zAQi`_ii4aSSt=64@M72Iy8_>MX3_`^-GZ}roIhx!S=rL#tlD2I_IP`G|HS>oqRyNM z`p%bt6`X$x8N|%$JjZljiL4;c?dT-Ja`>_l0fu4T=R&nF$~-sAlI5MJ3jUnk{XN60 z8_+g1I=a36;(h$>cS7>yA>$IdQHwO93^*!!f9r$&P7s_i56L0#G%`&-)m&E>YF&Zv QsF15;Vdnn?uAwLY2aGR0YybcN diff --git a/src/main.cpp b/src/main.cpp index 2c97a40..7d63754 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -150,7 +150,7 @@ int main(int argc, char* argv[]) { printf("** SCAN PERFORMANCE **\n"); printf("*****************************\n"); - for (int s = 4; s < 25; s++){ + for (int s = 4; s < 19; s++){ int ssize = 1 << s; int *u, *v; u = (int *)malloc(ssize*sizeof(int)); @@ -162,7 +162,7 @@ int main(int argc, char* argv[]) { auto start = std::chrono::high_resolution_clock::now(); StreamCompaction::CPU::scan(ssize, v, u); auto end = std::chrono::high_resolution_clock::now(); - double diff = std::chrono::duration_cast(end-start).count(); + double diff = std::chrono::duration_cast(end-start).count(); printf("CPU scan: %f\n", diff); zeroArray(ssize, v); diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu index 6ca121a..79edaec 100644 --- a/stream_compaction/efficient.cu +++ b/stream_compaction/efficient.cu @@ -66,7 +66,7 @@ void scan(int n, int *odata, const int *idata) { // Scan cudaEventRecord(start); for (int d = 0; d < ilog2ceil(m); d++){ - scanUp << <2, (int)m/2>> >(d, dev_pidata); + scanUp << > >(d, dev_pidata); } cudaEventRecord(stop); cudaEventSynchronize(stop); @@ -79,7 +79,7 @@ void scan(int n, int *odata, const int *idata) { cudaEventRecord(start); for (int d = ilog2ceil(m)-1; d >=0; d--){ - scanDown<<<2, (int)m/2>>>(d, dev_pidata); + scanDown<<>>(d, dev_pidata); } cudaEventRecord(stop); cudaEventSynchronize(stop); diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu index 310a799..a33d29c 100644 --- a/stream_compaction/naive.cu +++ b/stream_compaction/naive.cu @@ -41,7 +41,7 @@ void scan(int n, int *odata, const int *idata) { // Scan cudaEventRecord(start); for (int d = 1; d <= ilog2ceil(m); d++){ - scanCol<<<4, m/4>>>(d, dev_pidata); + scanCol<<>>(d, dev_pidata); } cudaEventRecord(stop); cudaEventSynchronize(stop); From e371ad0e822d1c55ea7ef7979ced949d191c8f6e Mon Sep 17 00:00:00 2001 From: Tongbo Sui Date: Sat, 12 Sep 2015 12:14:07 -0400 Subject: [PATCH 8/8] Update README.md --- README.md | 19 ------------------- 1 file changed, 19 deletions(-) diff --git a/README.md b/README.md index 384a224..da64423 100644 --- a/README.md +++ b/README.md @@ -34,10 +34,6 @@ All 4 implementations are ran on various array sizes from `2^4` to `2^18`. Due t ![](images/no-cpu-performance.jpg) - * Naive vs. Work-efficient - -![](images/custom-only-performance.jpg) - * Observations * Theoretically work-efficient scan should be faster than a naive scan. However somehow this is not the case in performance testing with small array sizes. A profile report shows the following which might contribute to the result: * Up-sweep is generating 100% more memory transactions than a naive scan. Adding down-sweep process, work-efficient method has about 150% more than naive method. @@ -213,18 +209,3 @@ Naive scan: 38.985985 Work-efficient scan: 108.208801 Thrust scan: 2.610464 ``` - -## Submit - -If you have modified any of the `CMakeLists.txt` files at all (aside from the -list of `SOURCE_FILES`), you must test that your project can build in Moore -100B/C. Beware of any build issues discussed on the Google Group. - -1. Open a GitHub pull request so that we can see that you have finished. - The title should be "Submission: YOUR NAME". -2. Send an email to the TA (gmail: kainino1+cis565@) with: - * **Subject**: in the form of `[CIS565] Project 2: PENNKEY` - * Direct link to your pull request on GitHub - * In the form of a grade (0-100+) with comments, evaluate your own - performance on the project. - * Feedback on the project itself, if any.