From 0366d4ba0628c863a67c38541e0f59b9841b9933 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Mon, 12 Sep 2022 17:24:42 -0400
Subject: [PATCH 01/17] Update README.md

---
 README.md | 10 +++++++---
 1 file changed, 7 insertions(+), 3 deletions(-)

diff --git a/README.md b/README.md
index 0e38ddb..ad1c0bd 100644
--- a/README.md
+++ b/README.md
@@ -3,9 +3,13 @@ CUDA Stream Compaction
 
 **University of Pennsylvania, CIS 565: GPU Programming and Architecture, Project 2**
 
-* (TODO) YOUR NAME HERE
-  * (TODO) [LinkedIn](), [personal website](), [twitter](), etc.
-* Tested on: (TODO) Windows 22, i7-2222 @ 2.22GHz 22GB, GTX 222 222MB (Moore 2222 Lab)
+Constance Wang
+  * [LinkedIn](https://www.linkedin.com/in/conswang/)
+
+Tested on AORUS 15P XD laptop with specs:  
+- Windows 11 22000.856  
+- 11th Gen Intel(R) Core(TM) i7-11800H @ 2.30GHz 2.30 GHz  
+- NVIDIA GeForce RTX 3070 Laptop GPU  
 
 ### (TODO: Your README)
 

From a8cfb893a2f309c2362075abf3a9811270d15363 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Fri, 16 Sep 2022 12:55:42 -0400
Subject: [PATCH 02/17] add cpu implementations

---
 stream_compaction/cpu.cu | 55 ++++++++++++++++++++++++++++++++++++++--
 1 file changed, 53 insertions(+), 2 deletions(-)

diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu
index 719fa11..cf65471 100644
--- a/stream_compaction/cpu.cu
+++ b/stream_compaction/cpu.cu
@@ -12,6 +12,17 @@ namespace StreamCompaction {
             return timer;
         }
 
+        void exclusiveScan(int n, int* odata, const int* idata) {
+          if (n <= 0) {
+            printf("Empty array!");
+            return;
+          }
+          odata[0] = 0; // identity
+          for (int i = 0; i < n - 1; ++i) {
+            odata[i + 1] = odata[i] + idata[i];
+          }
+        }
+
         /**
          * CPU scan (prefix sum).
          * For performance analysis, this is supposed to be a simple for loop.
@@ -20,6 +31,7 @@ namespace StreamCompaction {
         void scan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            exclusiveScan(n, odata, idata);
             timer().endCpuTimer();
         }
 
@@ -31,8 +43,15 @@ namespace StreamCompaction {
         int compactWithoutScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            int outputIndex = 0;
+            for (int inputIndex = 0; inputIndex < n; ++inputIndex) {
+              if (idata[inputIndex] != 0) {
+                odata[outputIndex] = idata[inputIndex];
+                outputIndex += 1;
+              }
+            }
             timer().endCpuTimer();
-            return -1;
+            return outputIndex;
         }
 
         /**
@@ -43,8 +62,40 @@ namespace StreamCompaction {
         int compactWithScan(int n, int *odata, const int *idata) {
             timer().startCpuTimer();
             // TODO
+            // calculate bool. array
+            int* boolData = new int[n];
+            if (boolData == NULL) {
+              printf("Failed to allocate boolData");
+              return -1;
+            }
+            for (int i = 0; i < n; ++i) {
+              boolData[i] = (idata[i] == 0) ? 0 : 1;
+            }
+
+            int* boolScan = new int[n];
+            if (boolScan == NULL) {
+              printf("Failed to allocate boolScan");
+              return -1;
+            }
+            exclusiveScan(n, boolScan, boolData);
+            //printf("boolScan ");
+            //for (int i = 0; i < n; ++i) {
+            //  printf("%d ", boolScan[i]);
+            //}
+
+            int outIndex = -1;
+            // Scatter
+            for (int i = 0; i < n; ++i) {
+              if (boolData[i] == 1) {
+                outIndex = boolScan[i];
+                odata[outIndex] = idata[i];
+              }
+            }
+            delete[] boolData;
+            delete[] boolScan;
             timer().endCpuTimer();
-            return -1;
+
+            return outIndex + 1;
         }
     }
 }

From dd07ee2470b88eefde0856d12e0b1c633b7f9be3 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Fri, 16 Sep 2022 15:53:36 -0400
Subject: [PATCH 03/17] working naive scan - inclusive, only for pow of 2

---
 stream_compaction/naive.cu | 54 ++++++++++++++++++++++++++++++++++++++
 1 file changed, 54 insertions(+)

diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu
index 4308876..57a006b 100644
--- a/stream_compaction/naive.cu
+++ b/stream_compaction/naive.cu
@@ -3,6 +3,12 @@
 #include "common.h"
 #include "naive.h"
 
+#include <device_launch_parameters.h>
+#include <device_functions.h>
+
+#define blockSize 128
+#define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__)
+
 namespace StreamCompaction {
     namespace Naive {
         using StreamCompaction::Common::PerformanceTimer;
@@ -11,7 +17,23 @@ namespace StreamCompaction {
             static PerformanceTimer timer;
             return timer;
         }
+        int* dev_idata;
+        int* dev_odata;
+
         // TODO: __global__
+        __global__ void kernScan(int n, int offset, int* odata, const int* idata) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= n) {
+            return;
+          }
+
+          if (index >= offset) {
+            odata[index] = idata[index - offset] + idata[index];
+          }
+          else {
+            odata[index] = idata[index];
+          }
+        }
 
         /**
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
@@ -19,6 +41,38 @@ namespace StreamCompaction {
         void scan(int n, int *odata, const int *idata) {
             timer().startGpuTimer();
             // TODO
+            cudaMalloc((void**)&dev_idata, n * sizeof(int));
+            cudaMalloc((void**)&dev_odata, n * sizeof(int));
+            checkCUDAErrorWithLine("cudaMalloc failed");
+
+            cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
+            checkCUDAErrorWithLine("memcpy idata failed!");
+
+            dim3 fullBlocksPerGrid((n - blockSize + 1) / blockSize);
+
+            int logOfN = ilog2ceil(n);
+
+            for (int d = 1; d <= logOfN; ++d) {
+              int offset = pow(2, d - 1);
+              kernScan << < fullBlocksPerGrid, blockSize >> > (n, offset, dev_odata, dev_idata);
+
+              // Needed to make sure we don't launch next iteration while prev is running
+              cudaDeviceSynchronize();
+
+              std::swap(dev_odata, dev_idata);
+            }
+
+            // Now result buffer is in dev_idata
+            cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
+            checkCUDAErrorWithLine("memcpy back to odata failed");
+
+            //printf("mewo mewomoemomemeow");
+            //for (int i = 0; i < n; ++i) {
+            //  printf("%d ", odata[i]);
+            //}
+
+            cudaFree(dev_idata);
+            cudaFree(dev_odata);
             timer().endGpuTimer();
         }
     }

From 84b3fac5e9593531950e176ddc356205a451be1c Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Fri, 16 Sep 2022 21:34:27 -0400
Subject: [PATCH 04/17] Working naive scan for non-powers of 2

---
 src/main.cpp               |  4 +--
 stream_compaction/naive.cu | 61 ++++++++++++++++++++++++++++++--------
 2 files changed, 51 insertions(+), 14 deletions(-)

diff --git a/src/main.cpp b/src/main.cpp
index 896ac2b..6ec45c8 100644
--- a/src/main.cpp
+++ b/src/main.cpp
@@ -13,8 +13,8 @@
 #include <stream_compaction/thrust.h>
 #include "testing_helpers.hpp"
 
-const int SIZE = 1 << 8; // feel free to change the size of array
-const int NPOT = SIZE - 3; // Non-Power-Of-Two
+const int SIZE = 1 << 8; // feel free to change the size of array = 256
+const int NPOT = SIZE - 3; // Non-Power-Of-Two = 253
 int *a = new int[SIZE];
 int *b = new int[SIZE];
 int *c = new int[SIZE];
diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu
index 57a006b..fae4e00 100644
--- a/stream_compaction/naive.cu
+++ b/stream_compaction/naive.cu
@@ -5,6 +5,7 @@
 
 #include <device_launch_parameters.h>
 #include <device_functions.h>
+#include <thrust/device_ptr.h>
 
 #define blockSize 128
 #define checkCUDAErrorWithLine(msg) checkCUDAError(msg, __LINE__)
@@ -19,6 +20,7 @@ namespace StreamCompaction {
         }
         int* dev_idata;
         int* dev_odata;
+        int* dev_unpadded_idata;
 
         // TODO: __global__
         __global__ void kernScan(int n, int offset, int* odata, const int* idata) {
@@ -35,26 +37,64 @@ namespace StreamCompaction {
           }
         }
 
+        // Pad the data with 1 zero at the beginning
+        // And enough zeroes at the end
+        // odata should be buffer of size paddedLength = the next power of two after and including (n + 1)
+        __global__ void kernShiftAndPadInput(int n, int paddedLength, int* odata, int* idata) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index == 0) {
+            odata[index] = 0;
+          }
+          else if (index <= n) {
+            odata[index] = idata[index - 1];
+          }
+          else if (index < paddedLength) {
+            odata[index] = 0;
+          }
+          else {
+            return;
+          }
+        }
+
+        void printCudaArray(int n, int* dev_array) {
+          int* tempArray = (int*) malloc(n * sizeof(int));
+          cudaMemcpy(tempArray, dev_array, n * sizeof(int), cudaMemcpyDeviceToHost);
+          printf("Print array -----------\n");
+          for (int i = 0; i < n; ++i) {
+            printf("%d ", tempArray[i]);
+          }
+          free(tempArray);
+        }
+
         /**
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
          */
         void scan(int n, int *odata, const int *idata) {
             timer().startGpuTimer();
             // TODO
-            cudaMalloc((void**)&dev_idata, n * sizeof(int));
-            cudaMalloc((void**)&dev_odata, n * sizeof(int));
+            
+            // Add 1 because we're going to offset by a zero for exclusive scan
+            int exponent = ilog2ceil(n + 1);
+            int paddedLength = pow(2, exponent);
+            // Input and output should be padded by 1s and 0s
+            cudaMalloc((void**)&dev_unpadded_idata, n * sizeof(int));
+            cudaMalloc((void**)&dev_idata, paddedLength * sizeof(int));
+            cudaMalloc((void**)&dev_odata, paddedLength * sizeof(int));
             checkCUDAErrorWithLine("cudaMalloc failed");
 
-            cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
+            cudaMemcpy(dev_unpadded_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
             checkCUDAErrorWithLine("memcpy idata failed!");
 
-            dim3 fullBlocksPerGrid((n - blockSize + 1) / blockSize);
+            dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
+
+            kernShiftAndPadInput<<<fullBlocksPerGrid, blockSize>>>
+              (n, paddedLength, dev_idata, dev_unpadded_idata);
 
-            int logOfN = ilog2ceil(n);
+            //printCudaArray(paddedLength, dev_idata);
 
-            for (int d = 1; d <= logOfN; ++d) {
+            for (int d = 1; d <= exponent; ++d) {
               int offset = pow(2, d - 1);
-              kernScan << < fullBlocksPerGrid, blockSize >> > (n, offset, dev_odata, dev_idata);
+              kernScan << < fullBlocksPerGrid, blockSize >> > (paddedLength, offset, dev_odata, dev_idata);
 
               // Needed to make sure we don't launch next iteration while prev is running
               cudaDeviceSynchronize();
@@ -63,16 +103,13 @@ namespace StreamCompaction {
             }
 
             // Now result buffer is in dev_idata
+            // We only need the first n elements of the total paddedLength elements
             cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
             checkCUDAErrorWithLine("memcpy back to odata failed");
 
-            //printf("mewo mewomoemomemeow");
-            //for (int i = 0; i < n; ++i) {
-            //  printf("%d ", odata[i]);
-            //}
-
             cudaFree(dev_idata);
             cudaFree(dev_odata);
+            cudaFree(dev_unpadded_idata);
             timer().endGpuTimer();
         }
     }

From d1b181e49b65dc9850777d0c3e46910cc1900ccb Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Fri, 16 Sep 2022 22:51:12 -0400
Subject: [PATCH 05/17] zero padding working without shift

---
 stream_compaction/common.cu    |  1 -
 stream_compaction/common.h     | 10 +++++++++
 stream_compaction/efficient.cu | 38 ++++++++++++++++++++++++++++++++++
 stream_compaction/naive.cu     | 10 ---------
 4 files changed, 48 insertions(+), 11 deletions(-)

diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu
index 2ed6d63..58bb3b1 100644
--- a/stream_compaction/common.cu
+++ b/stream_compaction/common.cu
@@ -17,7 +17,6 @@ void checkCUDAErrorFn(const char *msg, const char *file, int line) {
 
 namespace StreamCompaction {
     namespace Common {
-
         /**
          * Maps an array to an array of 0s and 1s for stream compaction. Elements
          * which map to 0 will be removed, and elements which map to 1 will be kept.
diff --git a/stream_compaction/common.h b/stream_compaction/common.h
index d2c1fed..5405eeb 100644
--- a/stream_compaction/common.h
+++ b/stream_compaction/common.h
@@ -30,6 +30,16 @@ inline int ilog2ceil(int x) {
     return x == 1 ? 0 : ilog2(x - 1) + 1;
 }
 
+inline void printCudaArray(int n, int* dev_array) {
+  int* tempArray = (int*)malloc(n * sizeof(int));
+  cudaMemcpy(tempArray, dev_array, n * sizeof(int), cudaMemcpyDeviceToHost);
+  printf("Print array -----------\n");
+  for (int i = 0; i < n; ++i) {
+    printf("%d ", tempArray[i]);
+  }
+  free(tempArray);
+}
+
 namespace StreamCompaction {
     namespace Common {
         __global__ void kernMapToBoolean(int n, int *bools, const int *idata);
diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu
index 2db346e..07de68c 100644
--- a/stream_compaction/efficient.cu
+++ b/stream_compaction/efficient.cu
@@ -3,6 +3,8 @@
 #include "common.h"
 #include "efficient.h"
 
+#define blockSize 128
+
 namespace StreamCompaction {
     namespace Efficient {
         using StreamCompaction::Common::PerformanceTimer;
@@ -12,12 +14,48 @@ namespace StreamCompaction {
             return timer;
         }
 
+        // unlike naive impl, this one doesn't shift the array
+        __global__ void kernPadArray(int n, int paddedLen, int* odata, int* idata) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index < n) {
+            odata[index] = idata[index];
+          }
+          else if (index < paddedLen) {
+            odata[index] = 0;
+          }
+        }
+
+        int* dev_unpadded_idata;
+        int* dev_idata;
+        int* dev_odata;
+
         /**
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
          */
         void scan(int n, int *odata, const int *idata) {
             timer().startGpuTimer();
             // TODO
+            // Pad array
+            int exponent = ilog2ceil(n);
+            int paddedLength = pow(2, exponent);
+            dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
+
+            cudaMalloc((void**)&dev_unpadded_idata, n * sizeof(int));
+            cudaMalloc((void**)&dev_idata, paddedLength * sizeof(int));
+            cudaMalloc((void**)&dev_odata, paddedLength * sizeof(int));
+
+            cudaMemcpy(dev_unpadded_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
+            checkCUDAErrorFn("Cuda memcpy idata no work");
+
+            kernPadArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_idata, dev_unpadded_idata);
+            printCudaArray(paddedLength, dev_idata);
+
+            // Build tree
+
+            cudaFree(dev_unpadded_idata);
+            cudaFree(dev_idata);
+            cudaFree(dev_odata);
+
             timer().endGpuTimer();
         }
 
diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu
index fae4e00..873aba8 100644
--- a/stream_compaction/naive.cu
+++ b/stream_compaction/naive.cu
@@ -56,16 +56,6 @@ namespace StreamCompaction {
           }
         }
 
-        void printCudaArray(int n, int* dev_array) {
-          int* tempArray = (int*) malloc(n * sizeof(int));
-          cudaMemcpy(tempArray, dev_array, n * sizeof(int), cudaMemcpyDeviceToHost);
-          printf("Print array -----------\n");
-          for (int i = 0; i < n; ++i) {
-            printf("%d ", tempArray[i]);
-          }
-          free(tempArray);
-        }
-
         /**
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
          */

From 7f786cf3ff88ac5adfa3133a2a43652001a18db4 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sat, 17 Sep 2022 19:52:59 -0400
Subject: [PATCH 06/17] working upsweep

---
 stream_compaction/efficient.cu | 34 ++++++++++++++++++++++++++++++----
 1 file changed, 30 insertions(+), 4 deletions(-)

diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu
index 07de68c..7ec0686 100644
--- a/stream_compaction/efficient.cu
+++ b/stream_compaction/efficient.cu
@@ -15,7 +15,7 @@ namespace StreamCompaction {
         }
 
         // unlike naive impl, this one doesn't shift the array
-        __global__ void kernPadArray(int n, int paddedLen, int* odata, int* idata) {
+        __global__ void kernPadArray(int n, int paddedLen, int* odata, const int* idata) {
           int index = threadIdx.x + (blockIdx.x * blockDim.x);
           if (index < n) {
             odata[index] = idata[index];
@@ -25,6 +25,21 @@ namespace StreamCompaction {
           }
         }
 
+        __global__ void kernUpsweep(int numThreads, int readStride, int* data) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= numThreads) {
+            return;
+          }
+
+          int writeStride = readStride * 2;
+
+          // Index of what element to write to is calculated using write stride
+          int writeIndex = (writeStride * index) + writeStride - 1;
+          int readIndex = (writeStride * index) + readStride - 1;
+
+          data[writeIndex] += data[readIndex];
+        }
+
         int* dev_unpadded_idata;
         int* dev_idata;
         int* dev_odata;
@@ -42,19 +57,30 @@ namespace StreamCompaction {
 
             cudaMalloc((void**)&dev_unpadded_idata, n * sizeof(int));
             cudaMalloc((void**)&dev_idata, paddedLength * sizeof(int));
-            cudaMalloc((void**)&dev_odata, paddedLength * sizeof(int));
 
             cudaMemcpy(dev_unpadded_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
             checkCUDAErrorFn("Cuda memcpy idata no work");
 
             kernPadArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_idata, dev_unpadded_idata);
+
+            // Build tree (upsweep)
+            // readStride = 2^depth, where depth goes from 0... log2n - 1... the stride between elements we read and sum
+            // writeStride = 2^(depth + 1)... the stride between indices of elements we store the sums in
+            for (int readStride = 1; readStride < paddedLength; readStride *= 2) {
+              int writeStride = readStride * 2;
+
+              int numThreads = paddedLength / writeStride; // one thread per element we write to
+              dim3 numBlocks((paddedLength + blockSize - 1) / blockSize);
+
+              kernUpsweep << <numBlocks, blockSize >> > (numThreads, readStride, dev_idata);
+            }
+
             printCudaArray(paddedLength, dev_idata);
 
-            // Build tree
+            cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
 
             cudaFree(dev_unpadded_idata);
             cudaFree(dev_idata);
-            cudaFree(dev_odata);
 
             timer().endGpuTimer();
         }

From 2449d0db6ccbcec8034a3d8eb081c8182a6e43a3 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sat, 17 Sep 2022 20:41:38 -0400
Subject: [PATCH 07/17] working work efficient scan

---
 stream_compaction/efficient.cu | 38 ++++++++++++++++++++++++++++++++--
 1 file changed, 36 insertions(+), 2 deletions(-)

diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu
index 7ec0686..61dd030 100644
--- a/stream_compaction/efficient.cu
+++ b/stream_compaction/efficient.cu
@@ -40,6 +40,22 @@ namespace StreamCompaction {
           data[writeIndex] += data[readIndex];
         }
 
+        __global__ void kernDownsweep(int numThreads, int writeStride, int* data) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= numThreads) {
+            return;
+          }
+
+          int readStride = writeStride * 2;
+
+          int leftChildIndex = index * readStride + writeStride - 1;
+          int rightChildIndex = index * readStride + readStride - 1; // right child is also where parent is stored
+
+          int temp = data[leftChildIndex];
+          data[leftChildIndex] = data[rightChildIndex];
+          data[rightChildIndex] += temp;
+        }
+
         int* dev_unpadded_idata;
         int* dev_idata;
         int* dev_odata;
@@ -70,12 +86,30 @@ namespace StreamCompaction {
               int writeStride = readStride * 2;
 
               int numThreads = paddedLength / writeStride; // one thread per element we write to
-              dim3 numBlocks((paddedLength + blockSize - 1) / blockSize);
+              dim3 numBlocks((numThreads + blockSize - 1) / blockSize);
 
               kernUpsweep << <numBlocks, blockSize >> > (numThreads, readStride, dev_idata);
+              cudaDeviceSynchronize();
             }
 
-            printCudaArray(paddedLength, dev_idata);
+            // Down sweep
+            // In down sweep, children now read info from parent. So writeStride = readStride / 2
+            // Write stride = n/2, n/4, ... 4, 2, 1, aka. 2^depth
+            // Read stride = 2^(depth + 1)
+            
+            // First set parent to 0
+            int zero = 0;
+            cudaMemcpy(dev_idata + paddedLength - 1, &zero, sizeof(int), cudaMemcpyHostToDevice);
+
+            for (int writeStride = paddedLength / 2; writeStride >= 1; writeStride = writeStride >> 1) {
+              int readStride = writeStride * 2;
+              
+              // now launch 1 thread per element we read from
+              int numThreads = paddedLength / readStride;
+              dim3 numBlocks((numThreads + blockSize - 1) / blockSize);
+
+              kernDownsweep << <numBlocks, blockSize >> > (numThreads, writeStride, dev_idata);
+            }
 
             cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
 

From dacca18ade85a5a5cb200fb895e65d8d53336502 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sat, 17 Sep 2022 21:29:41 -0400
Subject: [PATCH 08/17] gpu stream compaction appears to work

---
 stream_compaction/efficient.cu | 140 ++++++++++++++++++++++++---------
 1 file changed, 104 insertions(+), 36 deletions(-)

diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu
index 61dd030..4fabf61 100644
--- a/stream_compaction/efficient.cu
+++ b/stream_compaction/efficient.cu
@@ -56,9 +56,39 @@ namespace StreamCompaction {
           data[rightChildIndex] += temp;
         }
 
-        int* dev_unpadded_idata;
-        int* dev_idata;
-        int* dev_odata;
+        void scanImpl(int paddedLength, int* dev_idata) {
+          // Build tree (upsweep)
+          // readStride = 2^depth, where depth goes from 0... log2n - 1... the stride between elements we read and sum
+          // writeStride = 2^(depth + 1)... the stride between indices of elements we store the sums in
+          for (int readStride = 1; readStride < paddedLength; readStride *= 2) {
+            int writeStride = readStride * 2;
+
+            int numThreads = paddedLength / writeStride; // one thread per element we write to
+            dim3 numBlocks((numThreads + blockSize - 1) / blockSize);
+
+            kernUpsweep << <numBlocks, blockSize >> > (numThreads, readStride, dev_idata);
+            cudaDeviceSynchronize();
+          }
+
+          // Down sweep
+          // In down sweep, children now read info from parent. So writeStride = readStride / 2
+          // Write stride = n/2, n/4, ... 4, 2, 1, aka. 2^depth
+          // Read stride = 2^(depth + 1)
+
+          // First set parent to 0
+          int zero = 0;
+          cudaMemcpy(dev_idata + paddedLength - 1, &zero, sizeof(int), cudaMemcpyHostToDevice);
+
+          for (int writeStride = paddedLength / 2; writeStride >= 1; writeStride = writeStride >> 1) {
+            int readStride = writeStride * 2;
+
+            // now launch 1 thread per element we read from
+            int numThreads = paddedLength / readStride;
+            dim3 numBlocks((numThreads + blockSize - 1) / blockSize);
+
+            kernDownsweep << <numBlocks, blockSize >> > (numThreads, writeStride, dev_idata);
+          }
+        }
 
         /**
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
@@ -67,6 +97,9 @@ namespace StreamCompaction {
             timer().startGpuTimer();
             // TODO
             // Pad array
+            int* dev_unpadded_idata;
+            int* dev_idata;
+
             int exponent = ilog2ceil(n);
             int paddedLength = pow(2, exponent);
             dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
@@ -79,38 +112,7 @@ namespace StreamCompaction {
 
             kernPadArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_idata, dev_unpadded_idata);
 
-            // Build tree (upsweep)
-            // readStride = 2^depth, where depth goes from 0... log2n - 1... the stride between elements we read and sum
-            // writeStride = 2^(depth + 1)... the stride between indices of elements we store the sums in
-            for (int readStride = 1; readStride < paddedLength; readStride *= 2) {
-              int writeStride = readStride * 2;
-
-              int numThreads = paddedLength / writeStride; // one thread per element we write to
-              dim3 numBlocks((numThreads + blockSize - 1) / blockSize);
-
-              kernUpsweep << <numBlocks, blockSize >> > (numThreads, readStride, dev_idata);
-              cudaDeviceSynchronize();
-            }
-
-            // Down sweep
-            // In down sweep, children now read info from parent. So writeStride = readStride / 2
-            // Write stride = n/2, n/4, ... 4, 2, 1, aka. 2^depth
-            // Read stride = 2^(depth + 1)
-            
-            // First set parent to 0
-            int zero = 0;
-            cudaMemcpy(dev_idata + paddedLength - 1, &zero, sizeof(int), cudaMemcpyHostToDevice);
-
-            for (int writeStride = paddedLength / 2; writeStride >= 1; writeStride = writeStride >> 1) {
-              int readStride = writeStride * 2;
-              
-              // now launch 1 thread per element we read from
-              int numThreads = paddedLength / readStride;
-              dim3 numBlocks((numThreads + blockSize - 1) / blockSize);
-
-              kernDownsweep << <numBlocks, blockSize >> > (numThreads, writeStride, dev_idata);
-            }
-
+            scanImpl(paddedLength, dev_idata);
             cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
 
             cudaFree(dev_unpadded_idata);
@@ -119,6 +121,28 @@ namespace StreamCompaction {
             timer().endGpuTimer();
         }
 
+        __global__ void kernGetPaddedBoolArray(int n, int paddedLength, int* odata, const int* idata) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index < n) {
+            odata[index] = idata[index] == 0 ? 0 : 1;
+          }
+          else if (index < paddedLength) {
+            odata[index] = 0;
+          }
+        }
+
+        __global__ void kernScatter(int n, int* odata, int* idata, int *boolData, int *boolScan) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= n) {
+            return;
+          }
+
+          if (boolData[index] == 1) {
+            int writeIndex = boolScan[index];
+            odata[writeIndex] = idata[index];
+          }
+        }
+
         /**
          * Performs stream compaction on idata, storing the result into odata.
          * All zeroes are discarded.
@@ -131,8 +155,52 @@ namespace StreamCompaction {
         int compact(int n, int *odata, const int *idata) {
             timer().startGpuTimer();
             // TODO
+            int* dev_boolData;
+            int* dev_boolScan;
+            int* dev_idata;
+            int* dev_odata;
+            int exponent = ilog2ceil(n);
+            int paddedLength = pow(2, exponent);
+
+            cudaMalloc((void**)&dev_boolData, paddedLength * sizeof(int));
+            cudaMalloc((void**)&dev_boolScan, paddedLength * sizeof(int));
+            cudaMalloc((void**)&dev_idata, n * sizeof(int));
+            cudaMalloc((void**)&dev_odata, n * sizeof(int)); // don't need padding for scatter step
+            // copy input data
+            cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
+
+            dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
+            kernGetPaddedBoolArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_boolData, dev_idata);
+            // copy dev_boolData since the scan implementation is destructive
+            cudaMemcpy(dev_boolScan, dev_boolData, paddedLength * sizeof(int), cudaMemcpyDeviceToDevice);
+            scanImpl(paddedLength, dev_boolScan);
+
+            dim3 fullBlocksPerGridNonPadded((n + blockSize - 1) / blockSize);
+            kernScatter << <fullBlocksPerGridNonPadded, blockSize >> >
+              (n, dev_odata, dev_idata, dev_boolData, dev_boolScan);
+
+            cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost);
+
+            // How to get the total # of elements compacted:
+            // Get last (index n - 1) element of dev_boolScan and last (index n - 1) element of dev_boolData
+            // if dev_boolData[n - 1] = 0, then it's the last element
+            // Otherwise it's the last element + 1
+            
+            int lastElementIsIncluded, lastBoolScanVal, resultLength;
+
+            cudaMemcpy(&lastElementIsIncluded, dev_boolData + n - 1, sizeof(int), cudaMemcpyDeviceToHost);
+            cudaMemcpy(&lastBoolScanVal, dev_boolScan + n - 1, sizeof(int), cudaMemcpyDeviceToHost);
+            
+            resultLength = lastElementIsIncluded == 1 ? lastBoolScanVal + 1 : lastBoolScanVal;
+
+            cudaFree(dev_boolData);
+            cudaFree(dev_boolScan);
+            cudaFree(dev_idata);
+            cudaFree(dev_odata);
+
             timer().endGpuTimer();
-            return -1;
+
+            return resultLength;
         }
     }
 }

From 303672fe02eeb210cb9c3fcee20a47a4e691135b Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sat, 17 Sep 2022 22:03:52 -0400
Subject: [PATCH 09/17] update gpu timers and add thrust scan

---
 stream_compaction/efficient.cu | 20 +++++++++++++-------
 stream_compaction/naive.cu     |  6 ++++--
 stream_compaction/thrust.cu    | 15 +++++++++++++++
 3 files changed, 32 insertions(+), 9 deletions(-)

diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu
index 4fabf61..025bf38 100644
--- a/stream_compaction/efficient.cu
+++ b/stream_compaction/efficient.cu
@@ -94,7 +94,6 @@ namespace StreamCompaction {
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
          */
         void scan(int n, int *odata, const int *idata) {
-            timer().startGpuTimer();
             // TODO
             // Pad array
             int* dev_unpadded_idata;
@@ -110,15 +109,20 @@ namespace StreamCompaction {
             cudaMemcpy(dev_unpadded_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
             checkCUDAErrorFn("Cuda memcpy idata no work");
 
+            // start timer after all initial memcpys
+            timer().startGpuTimer();
+
             kernPadArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_idata, dev_unpadded_idata);
 
             scanImpl(paddedLength, dev_idata);
+
+            // end timer before copying result array
+            timer().endGpuTimer();
+
             cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
 
             cudaFree(dev_unpadded_idata);
             cudaFree(dev_idata);
-
-            timer().endGpuTimer();
         }
 
         __global__ void kernGetPaddedBoolArray(int n, int paddedLength, int* odata, const int* idata) {
@@ -153,7 +157,6 @@ namespace StreamCompaction {
          * @returns      The number of elements remaining after compaction.
          */
         int compact(int n, int *odata, const int *idata) {
-            timer().startGpuTimer();
             // TODO
             int* dev_boolData;
             int* dev_boolScan;
@@ -166,9 +169,11 @@ namespace StreamCompaction {
             cudaMalloc((void**)&dev_boolScan, paddedLength * sizeof(int));
             cudaMalloc((void**)&dev_idata, n * sizeof(int));
             cudaMalloc((void**)&dev_odata, n * sizeof(int)); // don't need padding for scatter step
-            // copy input data
             cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
 
+            // Start timer after copying input data
+            timer().startGpuTimer();
+
             dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
             kernGetPaddedBoolArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_boolData, dev_idata);
             // copy dev_boolData since the scan implementation is destructive
@@ -179,6 +184,9 @@ namespace StreamCompaction {
             kernScatter << <fullBlocksPerGridNonPadded, blockSize >> >
               (n, dev_odata, dev_idata, dev_boolData, dev_boolScan);
 
+            // End timer before copying odata out
+            timer().endGpuTimer();
+
             cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost);
 
             // How to get the total # of elements compacted:
@@ -198,8 +206,6 @@ namespace StreamCompaction {
             cudaFree(dev_idata);
             cudaFree(dev_odata);
 
-            timer().endGpuTimer();
-
             return resultLength;
         }
     }
diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu
index 873aba8..a9da4c8 100644
--- a/stream_compaction/naive.cu
+++ b/stream_compaction/naive.cu
@@ -60,7 +60,6 @@ namespace StreamCompaction {
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
          */
         void scan(int n, int *odata, const int *idata) {
-            timer().startGpuTimer();
             // TODO
             
             // Add 1 because we're going to offset by a zero for exclusive scan
@@ -75,6 +74,8 @@ namespace StreamCompaction {
             cudaMemcpy(dev_unpadded_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
             checkCUDAErrorWithLine("memcpy idata failed!");
 
+            timer().startGpuTimer();
+
             dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
 
             kernShiftAndPadInput<<<fullBlocksPerGrid, blockSize>>>
@@ -92,6 +93,8 @@ namespace StreamCompaction {
               std::swap(dev_odata, dev_idata);
             }
 
+            timer().endGpuTimer();
+
             // Now result buffer is in dev_idata
             // We only need the first n elements of the total paddedLength elements
             cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
@@ -100,7 +103,6 @@ namespace StreamCompaction {
             cudaFree(dev_idata);
             cudaFree(dev_odata);
             cudaFree(dev_unpadded_idata);
-            timer().endGpuTimer();
         }
     }
 }
diff --git a/stream_compaction/thrust.cu b/stream_compaction/thrust.cu
index 1def45e..27831c0 100644
--- a/stream_compaction/thrust.cu
+++ b/stream_compaction/thrust.cu
@@ -18,11 +18,26 @@ namespace StreamCompaction {
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
          */
         void scan(int n, int *odata, const int *idata) {
+            int* dev_idata;
+            int* dev_odata;
+            cudaMalloc((void**)&dev_idata, n * sizeof(int));
+            cudaMalloc((void**)&dev_odata, n * sizeof(int));
+
+            cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
+            thrust::device_ptr<int> thrust_dev_idata(dev_idata);
+            thrust::device_ptr<int> thrust_dev_odata(dev_odata);
+            
+            // Start timer after input data is copied
             timer().startGpuTimer();
             // TODO use `thrust::exclusive_scan`
             // example: for device_vectors dv_in and dv_out:
             // thrust::exclusive_scan(dv_in.begin(), dv_in.end(), dv_out.begin());
+            thrust::exclusive_scan(thrust_dev_idata, thrust_dev_idata + n, thrust_dev_odata);
+
             timer().endGpuTimer();
+
+            // Copy the data back (not included in timing)
+            cudaMemcpy(odata, dev_odata, n * sizeof(int), cudaMemcpyDeviceToHost);
         }
     }
 }

From afc875f44e464141810d34e37ad06cc294b012f4 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 19:51:02 -0400
Subject: [PATCH 10/17] working radish sort

---
 performance-at-last-commit.txt | 55 +++++++++++++++++++
 src/main.cpp                   | 37 ++++++++++++-
 stream_compaction/common.cu    | 20 +++++++
 stream_compaction/common.h     |  4 ++
 stream_compaction/cpu.cu       | 44 ++++++++++++++++
 stream_compaction/cpu.h        |  2 +
 stream_compaction/efficient.cu | 25 +--------
 stream_compaction/efficient.h  |  2 +
 stream_compaction/naive.cu     | 96 ++++++++++++++++++++++++++++++++++
 stream_compaction/naive.h      |  1 +
 10 files changed, 262 insertions(+), 24 deletions(-)
 create mode 100644 performance-at-last-commit.txt

diff --git a/performance-at-last-commit.txt b/performance-at-last-commit.txt
new file mode 100644
index 0000000..b1c619c
--- /dev/null
+++ b/performance-at-last-commit.txt
@@ -0,0 +1,55 @@
+On length = 1 << 18
+
+****************
+** SCAN TESTS **
+****************
+    [   4  25  22   2  14  31  30   0  44  13  11  45  24 ...   5   0 ]
+==== cpu scan, power-of-two ====
+   elapsed time: 0.3819ms    (std::chrono Measured)
+    [   0   4  29  51  53  67  98 128 128 172 185 196 241 ... 6433433 6433438 ]
+==== cpu scan, non-power-of-two ====
+   elapsed time: 0.4017ms    (std::chrono Measured)
+    [   0   4  29  51  53  67  98 128 128 172 185 196 241 ... 6433393 6433393 ]
+    passed
+==== naive scan, power-of-two ====
+   elapsed time: 0.327168ms    (CUDA Measured)
+    passed
+==== naive scan, non-power-of-two ====
+   elapsed time: 0.229856ms    (CUDA Measured)
+    passed
+==== work-efficient scan, power-of-two ====
+   elapsed time: 0.331456ms    (CUDA Measured)
+    passed
+==== work-efficient scan, non-power-of-two ====
+   elapsed time: 0.343808ms    (CUDA Measured)
+    passed
+==== thrust scan, power-of-two ====
+   elapsed time: 0.146592ms    (CUDA Measured)
+    passed
+==== thrust scan, non-power-of-two ====
+   elapsed time: 0.190304ms    (CUDA Measured)
+    passed
+
+*****************************
+** STREAM COMPACTION TESTS **
+*****************************
+    [   3   0   0   1   3   3   3   3   1   1   1   1   2 ...   1   0 ]
+==== cpu compact without scan, power-of-two ====
+   elapsed time: 0.5161ms    (std::chrono Measured)
+    [   3   1   3   3   3   3   1   1   1   1   2   3   3 ...   2   1 ]
+    passed
+==== cpu compact without scan, non-power-of-two ====
+   elapsed time: 0.5447ms    (std::chrono Measured)
+    [   3   1   3   3   3   3   1   1   1   1   2   3   3 ...   2   3 ]
+    passed
+==== cpu compact with scan ====
+   elapsed time: 1.4394ms    (std::chrono Measured)
+    [   3   1   3   3   3   3   1   1   1   1   2   3   3 ...   2   1 ]
+    passed
+==== work-efficient compact, power-of-two ====
+   elapsed time: 0.3688ms    (CUDA Measured)
+    passed
+==== work-efficient compact, non-power-of-two ====
+   elapsed time: 0.500928ms    (CUDA Measured)
+    passed
+Appuyez sur une touche pour continuer...
\ No newline at end of file
diff --git a/src/main.cpp b/src/main.cpp
index 6ec45c8..1abbd00 100644
--- a/src/main.cpp
+++ b/src/main.cpp
@@ -13,7 +13,7 @@
 #include <stream_compaction/thrust.h>
 #include "testing_helpers.hpp"
 
-const int SIZE = 1 << 8; // feel free to change the size of array = 256
+const int SIZE = 1 << 15; // feel free to change the size of array = 256
 const int NPOT = SIZE - 3; // Non-Power-Of-Two = 253
 int *a = new int[SIZE];
 int *b = new int[SIZE];
@@ -147,6 +147,41 @@ int main(int argc, char* argv[]) {
     //printArray(count, c, true);
     printCmpLenResult(count, expectedNPOT, b, c);
 
+    printf("\n");
+    printf("*****************************\n");
+    printf("** RADIX SORT TESTS **\n");
+    printf("*****************************\n");
+
+    genArray(SIZE - 1, a, 4);  // Leave a 0 at the end to test that edge case
+    a[SIZE - 1] = 0;
+    printArray(SIZE, a, true);
+
+    #define RADIX_NUM_BITS ilog2ceil(4)
+
+    zeroArray(SIZE, b);
+    printDesc("cpu radix sort, power-of-two");
+    StreamCompaction::CPU::radixSort(SIZE, RADIX_NUM_BITS, b, a);
+    //printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
+    printArray(SIZE, b, true);
+
+    zeroArray(SIZE, c);
+    printDesc("radix sort, power of two");
+    StreamCompaction::Naive::radixSort(SIZE, RADIX_NUM_BITS, c, a);
+    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
+    printCmpResult(SIZE, b, c); // TODO: add cpu impl and write to b for comparison
+
+    zeroArray(SIZE, b);
+    printDesc("cpu radix sort, non-power-of-two");
+    StreamCompaction::CPU::radixSort(NPOT, RADIX_NUM_BITS, b, a);
+    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
+    printArray(SIZE, b, true);
+
+    zeroArray(SIZE, c);
+    printDesc("radix sort, non-power of two");
+    StreamCompaction::Naive::radixSort(NPOT, RADIX_NUM_BITS, c, a);
+    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
+    printCmpResult(NPOT, b, c); // TODO: add cpu impl and write to b for comparison
+
     system("pause"); // stop Win32 console from closing on exit
     delete[] a;
     delete[] b;
diff --git a/stream_compaction/common.cu b/stream_compaction/common.cu
index 58bb3b1..37f6a8f 100644
--- a/stream_compaction/common.cu
+++ b/stream_compaction/common.cu
@@ -34,5 +34,25 @@ namespace StreamCompaction {
             // TODO
         }
 
+        // unlike naive impl, this one doesn't shift the array
+        __global__ void kernPadArray(int n, int paddedLen, int* odata, const int* idata) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index < n) {
+            odata[index] = idata[index];
+          }
+          else if (index < paddedLen) {
+            odata[index] = 0;
+          }
+        }
+
+        __global__ void kernGetPaddedBoolArray(int n, int paddedLength, int* odata, const int* idata) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index < n) {
+            odata[index] = idata[index] == 0 ? 0 : 1;
+          }
+          else if (index < paddedLength) {
+            odata[index] = 0;
+          }
+        }
     }
 }
diff --git a/stream_compaction/common.h b/stream_compaction/common.h
index 5405eeb..575ce35 100644
--- a/stream_compaction/common.h
+++ b/stream_compaction/common.h
@@ -47,6 +47,10 @@ namespace StreamCompaction {
         __global__ void kernScatter(int n, int *odata,
                 const int *idata, const int *bools, const int *indices);
 
+        __global__ void kernPadArray(int n, int paddedLength, int* odata, const int* idata);
+
+        __global__ void kernGetPaddedBoolArray(int n, int paddedLength, int* odata, const int* idata);
+
         /**
         * This class is used for timing the performance
         * Uncopyable and unmovable
diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu
index cf65471..eee979c 100644
--- a/stream_compaction/cpu.cu
+++ b/stream_compaction/cpu.cu
@@ -97,5 +97,49 @@ namespace StreamCompaction {
 
             return outIndex + 1;
         }
+
+        void split(int n, int bitmask, int* odata, const int* idata) {
+          // write 0s to odata in 1 pass, 1s in the next pass
+          int zeroCount = 0;
+          for (int i = 0; i < n; ++i) {
+            if ((idata[i] & bitmask) == 0) {
+              odata[zeroCount] = idata[i];
+              ++zeroCount; // eg. we write a total of 123 elts with 0s, then indices 0...122s are filled
+            }
+          }
+
+          int oneCount = 0;
+          for (int i = 0; i < n; ++i) {
+            if ((idata[i] & bitmask) != 0) {
+              odata[zeroCount + oneCount] = idata[i];
+              ++oneCount;
+            }
+          }
+        }
+
+        // CPU radix sort implementation
+        void radixSort(int n, int numBits, int* odata, const int* idata)
+        {
+          // use temporary buffers so we don't destroy idata
+          int* temp_odata = (int*) malloc(n * sizeof(int));
+          int* temp_odata2 = (int*)malloc(n * sizeof(int));
+          if (temp_odata == NULL || temp_odata2 == NULL) {
+            printf("ERROR - failed to allocate temp_odata");
+            return;
+          }
+          memcpy(temp_odata, idata, n * sizeof(int));
+
+          int bitmask = 1;
+          for (int i = 0; i < numBits; ++i) {
+            split(n, bitmask, temp_odata2, temp_odata);
+            std::swap(temp_odata2, temp_odata); // at end of loop, temp_odata always has most updated info
+            bitmask <<= 1;
+          }
+
+          memcpy(odata, temp_odata, n * sizeof(int));
+
+          free(temp_odata);
+          free(temp_odata2);
+        }
     }
 }
diff --git a/stream_compaction/cpu.h b/stream_compaction/cpu.h
index 873c047..9d9e8cd 100644
--- a/stream_compaction/cpu.h
+++ b/stream_compaction/cpu.h
@@ -11,5 +11,7 @@ namespace StreamCompaction {
         int compactWithoutScan(int n, int *odata, const int *idata);
 
         int compactWithScan(int n, int *odata, const int *idata);
+
+        void radixSort(int n, int numBits, int* odata, const int* idata);
     }
 }
diff --git a/stream_compaction/efficient.cu b/stream_compaction/efficient.cu
index 025bf38..a6d6010 100644
--- a/stream_compaction/efficient.cu
+++ b/stream_compaction/efficient.cu
@@ -14,17 +14,6 @@ namespace StreamCompaction {
             return timer;
         }
 
-        // unlike naive impl, this one doesn't shift the array
-        __global__ void kernPadArray(int n, int paddedLen, int* odata, const int* idata) {
-          int index = threadIdx.x + (blockIdx.x * blockDim.x);
-          if (index < n) {
-            odata[index] = idata[index];
-          }
-          else if (index < paddedLen) {
-            odata[index] = 0;
-          }
-        }
-
         __global__ void kernUpsweep(int numThreads, int readStride, int* data) {
           int index = threadIdx.x + (blockIdx.x * blockDim.x);
           if (index >= numThreads) {
@@ -112,7 +101,7 @@ namespace StreamCompaction {
             // start timer after all initial memcpys
             timer().startGpuTimer();
 
-            kernPadArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_idata, dev_unpadded_idata);
+            Common::kernPadArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_idata, dev_unpadded_idata);
 
             scanImpl(paddedLength, dev_idata);
 
@@ -125,16 +114,6 @@ namespace StreamCompaction {
             cudaFree(dev_idata);
         }
 
-        __global__ void kernGetPaddedBoolArray(int n, int paddedLength, int* odata, const int* idata) {
-          int index = threadIdx.x + (blockIdx.x * blockDim.x);
-          if (index < n) {
-            odata[index] = idata[index] == 0 ? 0 : 1;
-          }
-          else if (index < paddedLength) {
-            odata[index] = 0;
-          }
-        }
-
         __global__ void kernScatter(int n, int* odata, int* idata, int *boolData, int *boolScan) {
           int index = threadIdx.x + (blockIdx.x * blockDim.x);
           if (index >= n) {
@@ -175,7 +154,7 @@ namespace StreamCompaction {
             timer().startGpuTimer();
 
             dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
-            kernGetPaddedBoolArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_boolData, dev_idata);
+            Common::kernGetPaddedBoolArray << <fullBlocksPerGrid, blockSize >> > (n, paddedLength, dev_boolData, dev_idata);
             // copy dev_boolData since the scan implementation is destructive
             cudaMemcpy(dev_boolScan, dev_boolData, paddedLength * sizeof(int), cudaMemcpyDeviceToDevice);
             scanImpl(paddedLength, dev_boolScan);
diff --git a/stream_compaction/efficient.h b/stream_compaction/efficient.h
index 803cb4f..b19f4d5 100644
--- a/stream_compaction/efficient.h
+++ b/stream_compaction/efficient.h
@@ -8,6 +8,8 @@ namespace StreamCompaction {
 
         void scan(int n, int *odata, const int *idata);
 
+        void scanImpl(int paddedLength, int* dev_data);
+
         int compact(int n, int *odata, const int *idata);
     }
 }
diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu
index a9da4c8..7e4ba53 100644
--- a/stream_compaction/naive.cu
+++ b/stream_compaction/naive.cu
@@ -2,6 +2,7 @@
 #include <cuda_runtime.h>
 #include "common.h"
 #include "naive.h"
+#include "efficient.h"
 
 #include <device_launch_parameters.h>
 #include <device_functions.h>
@@ -104,5 +105,100 @@ namespace StreamCompaction {
             cudaFree(dev_odata);
             cudaFree(dev_unpadded_idata);
         }
+
+        // oTrueData[index] is 1 if idata[index] has bit 1 at position i
+        // oFalseData is just inverted oTrueData
+        // bit mask is of the form (in binary) 000...1...00
+        __global__ void kernGetPaddedBoolArray
+          (int n, int paddedLength, int bitMask, int* falsesData, int* idata) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index < n) {
+            falsesData[index] = ((bitMask & idata[index]) == 0) ? 1 : 0;
+          }
+          else if (index < paddedLength) {
+            // Since we pad at the end, we need to pad end elements with bit = 1 (aka. falses = 0)
+            falsesData[index] = 0;
+          }
+        }
+
+        __global__ void kernGetScatterAddresses(int n, int *odata, const int* falsesData, const int *falsesScanData, int totalFalses) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= n) {
+            return;
+          }
+          odata[index] = falsesData[index] == 1 
+            ? falsesScanData[index]
+            : index - falsesScanData[index] + totalFalses; // the "true" scan data offset by totalFalses
+        }
+
+        __global__ void kernScatter(int n, int* odata, const int* idata, const int* scatterAddresses) {
+          int index = threadIdx.x + (blockIdx.x * blockDim.x);
+          if (index >= n) {
+            return;
+          }
+          odata[scatterAddresses[index]] = idata[index];
+        }
+
+        // Let's not worry about 2's complement for now...
+        // Assume we only have positive integers
+        void radixSort(int n, int numBits, int* odata, const int* idata) {
+          int* dev_idata, * dev_odata; // two buffers needed because scatter step has race conditions
+          int *dev_falsesData, * dev_falsesScanData, * dev_scatterAddresses;
+
+          int exponent = ilog2ceil(n);
+          int paddedLength = pow(2, exponent);
+          dim3 fullBlocksPerGridUnpadded((n + blockSize - 1) / blockSize);
+          dim3 fullBlocksPerGridPadded((paddedLength + blockSize - 1) / blockSize);
+
+          cudaMalloc((void**)&dev_idata, n * sizeof(int));
+          cudaMemcpy(dev_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
+
+          cudaMalloc((void**)&dev_odata, n * sizeof(int));
+
+          cudaMalloc((void**)&dev_falsesData, paddedLength * sizeof(int));
+          cudaMalloc((void**)&dev_falsesScanData, paddedLength * sizeof(int));
+          cudaMalloc((void**)&dev_scatterAddresses, n * sizeof(int)); 
+
+          int bitmask = 1;
+          for (int i = 0; i < numBits; ++i) {
+            // need indices from 0... paddedLength - 1
+            kernGetPaddedBoolArray << < fullBlocksPerGridPadded, blockSize >> >
+              (n, paddedLength, bitmask, dev_falsesData, dev_idata);
+            
+            cudaMemcpy(dev_falsesScanData, dev_falsesData, paddedLength * sizeof(int), cudaMemcpyDeviceToDevice);
+            Efficient::scanImpl(paddedLength, dev_falsesScanData);
+
+            // Calculate total falses = dev_falsesScanData[n-1] + dev_falsesData[n-1]
+            // TODO: just make it a function
+            int lastElementIsIncluded, lastBoolScanVal, totalFalses;
+            cudaMemcpy(&lastElementIsIncluded, dev_falsesData + n - 1, sizeof(int), cudaMemcpyDeviceToHost);
+            cudaMemcpy(&lastBoolScanVal, dev_falsesScanData + n - 1, sizeof(int), cudaMemcpyDeviceToHost);
+            totalFalses = lastElementIsIncluded + lastBoolScanVal;
+            
+            kernGetScatterAddresses << <fullBlocksPerGridUnpadded, blockSize >> >
+              (n, dev_scatterAddresses, dev_falsesData, dev_falsesScanData, totalFalses);
+
+            kernScatter<<<fullBlocksPerGridUnpadded, blockSize>>>(n, dev_odata, dev_idata, dev_scatterAddresses);
+
+            cudaDeviceSynchronize();
+
+            //printf("MEOW total falses = %d\n", totalFalses);
+            //printCudaArray(n, dev_falsesData);
+            //printCudaArray(n, dev_falsesScanData);
+            //printCudaArray(n, dev_scatterAddresses);
+            //printCudaArray(n, dev_odata);
+
+            bitmask = bitmask << 1; // eg. ...0010 => ...0100
+            std::swap(dev_odata, dev_idata); // dev_idata always has latest
+          }
+
+          cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
+
+          cudaFree(dev_idata);
+          cudaFree(dev_odata);
+          cudaFree(dev_falsesData);
+          cudaFree(dev_falsesScanData);
+          cudaFree(dev_scatterAddresses);
+        }
     }
 }
diff --git a/stream_compaction/naive.h b/stream_compaction/naive.h
index 37dcb06..86c4303 100644
--- a/stream_compaction/naive.h
+++ b/stream_compaction/naive.h
@@ -7,5 +7,6 @@ namespace StreamCompaction {
         StreamCompaction::Common::PerformanceTimer& timer();
 
         void scan(int n, int *odata, const int *idata);
+        void radixSort(int n, int numBits, int* odata, const int* idata);
     }
 }

From 34adb35e6a64e61c5ddf5113c0bfb3c31a5b0a1a Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 19:54:16 -0400
Subject: [PATCH 11/17] add timers

---
 stream_compaction/cpu.cu   | 4 ++++
 stream_compaction/naive.cu | 4 ++++
 2 files changed, 8 insertions(+)

diff --git a/stream_compaction/cpu.cu b/stream_compaction/cpu.cu
index eee979c..7196043 100644
--- a/stream_compaction/cpu.cu
+++ b/stream_compaction/cpu.cu
@@ -129,6 +129,8 @@ namespace StreamCompaction {
           }
           memcpy(temp_odata, idata, n * sizeof(int));
 
+          timer().startCpuTimer();
+
           int bitmask = 1;
           for (int i = 0; i < numBits; ++i) {
             split(n, bitmask, temp_odata2, temp_odata);
@@ -136,6 +138,8 @@ namespace StreamCompaction {
             bitmask <<= 1;
           }
 
+          timer().endCpuTimer(); // end before final memcpy (which just transfers output to odata)
+
           memcpy(odata, temp_odata, n * sizeof(int));
 
           free(temp_odata);
diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu
index 7e4ba53..7214625 100644
--- a/stream_compaction/naive.cu
+++ b/stream_compaction/naive.cu
@@ -159,6 +159,8 @@ namespace StreamCompaction {
           cudaMalloc((void**)&dev_falsesScanData, paddedLength * sizeof(int));
           cudaMalloc((void**)&dev_scatterAddresses, n * sizeof(int)); 
 
+          timer().startGpuTimer();
+
           int bitmask = 1;
           for (int i = 0; i < numBits; ++i) {
             // need indices from 0... paddedLength - 1
@@ -192,6 +194,8 @@ namespace StreamCompaction {
             std::swap(dev_odata, dev_idata); // dev_idata always has latest
           }
 
+          timer().endGpuTimer();
+
           cudaMemcpy(odata, dev_idata, n * sizeof(int), cudaMemcpyDeviceToHost);
 
           cudaFree(dev_idata);

From 802fc2997a68ab78fffbf643b7853b365c2797a4 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 21:05:38 -0400
Subject: [PATCH 12/17] remove unneeded padding from naive scan, fix test
 printing wrong number of elts

---
 README.md                  |  9 ++++-----
 src/main.cpp               |  8 ++++----
 stream_compaction/naive.cu | 41 ++++++++------------------------------
 3 files changed, 16 insertions(+), 42 deletions(-)

diff --git a/README.md b/README.md
index ad1c0bd..a4aff72 100644
--- a/README.md
+++ b/README.md
@@ -11,8 +11,7 @@ Tested on AORUS 15P XD laptop with specs:
 - 11th Gen Intel(R) Core(TM) i7-11800H @ 2.30GHz 2.30 GHz  
 - NVIDIA GeForce RTX 3070 Laptop GPU  
 
-### (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.)
-
+I implemented the following parallel algorithms on the GPU and benchmarked them against my own CPU implementations and Thrust on the GPU:  
+- Scan
+- Stream compaction
+- Radix sort
\ No newline at end of file
diff --git a/src/main.cpp b/src/main.cpp
index 1abbd00..0abb168 100644
--- a/src/main.cpp
+++ b/src/main.cpp
@@ -13,7 +13,7 @@
 #include <stream_compaction/thrust.h>
 #include "testing_helpers.hpp"
 
-const int SIZE = 1 << 15; // feel free to change the size of array = 256
+const int SIZE = 1 << 12; // feel free to change the size of array = 256
 const int NPOT = SIZE - 3; // Non-Power-Of-Two = 253
 int *a = new int[SIZE];
 int *b = new int[SIZE];
@@ -152,11 +152,11 @@ int main(int argc, char* argv[]) {
     printf("** RADIX SORT TESTS **\n");
     printf("*****************************\n");
 
-    genArray(SIZE - 1, a, 4);  // Leave a 0 at the end to test that edge case
+    genArray(SIZE - 1, a, 696969);  // Leave a 0 at the end to test that edge case
     a[SIZE - 1] = 0;
     printArray(SIZE, a, true);
 
-    #define RADIX_NUM_BITS ilog2ceil(4)
+    #define RADIX_NUM_BITS ilog2ceil(696969)
 
     zeroArray(SIZE, b);
     printDesc("cpu radix sort, power-of-two");
@@ -174,7 +174,7 @@ int main(int argc, char* argv[]) {
     printDesc("cpu radix sort, non-power-of-two");
     StreamCompaction::CPU::radixSort(NPOT, RADIX_NUM_BITS, b, a);
     printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
-    printArray(SIZE, b, true);
+    printArray(NPOT, b, true);
 
     zeroArray(SIZE, c);
     printDesc("radix sort, non-power of two");
diff --git a/stream_compaction/naive.cu b/stream_compaction/naive.cu
index 7214625..768b02c 100644
--- a/stream_compaction/naive.cu
+++ b/stream_compaction/naive.cu
@@ -19,9 +19,6 @@ namespace StreamCompaction {
             static PerformanceTimer timer;
             return timer;
         }
-        int* dev_idata;
-        int* dev_odata;
-        int* dev_unpadded_idata;
 
         // TODO: __global__
         __global__ void kernScan(int n, int offset, int* odata, const int* idata) {
@@ -38,52 +35,31 @@ namespace StreamCompaction {
           }
         }
 
-        // Pad the data with 1 zero at the beginning
-        // And enough zeroes at the end
-        // odata should be buffer of size paddedLength = the next power of two after and including (n + 1)
-        __global__ void kernShiftAndPadInput(int n, int paddedLength, int* odata, int* idata) {
-          int index = threadIdx.x + (blockIdx.x * blockDim.x);
-          if (index == 0) {
-            odata[index] = 0;
-          }
-          else if (index <= n) {
-            odata[index] = idata[index - 1];
-          }
-          else if (index < paddedLength) {
-            odata[index] = 0;
-          }
-          else {
-            return;
-          }
-        }
-
         /**
          * Performs prefix-sum (aka scan) on idata, storing the result into odata.
          */
         void scan(int n, int *odata, const int *idata) {
             // TODO
-            
+            int* dev_idata;
+            int* dev_odata;
+
             // Add 1 because we're going to offset by a zero for exclusive scan
+            int paddedLength = n + 1;
             int exponent = ilog2ceil(n + 1);
-            int paddedLength = pow(2, exponent);
-            // Input and output should be padded by 1s and 0s
-            cudaMalloc((void**)&dev_unpadded_idata, n * sizeof(int));
+
             cudaMalloc((void**)&dev_idata, paddedLength * sizeof(int));
             cudaMalloc((void**)&dev_odata, paddedLength * sizeof(int));
             checkCUDAErrorWithLine("cudaMalloc failed");
 
-            cudaMemcpy(dev_unpadded_idata, idata, n * sizeof(int), cudaMemcpyHostToDevice);
+            int identity = 0;
+            cudaMemcpy(dev_idata, &identity, sizeof(int), cudaMemcpyHostToDevice);
+            cudaMemcpy(dev_idata + 1, idata, n * sizeof(int), cudaMemcpyHostToDevice);
             checkCUDAErrorWithLine("memcpy idata failed!");
 
             timer().startGpuTimer();
 
             dim3 fullBlocksPerGrid((paddedLength + blockSize - 1) / blockSize);
 
-            kernShiftAndPadInput<<<fullBlocksPerGrid, blockSize>>>
-              (n, paddedLength, dev_idata, dev_unpadded_idata);
-
-            //printCudaArray(paddedLength, dev_idata);
-
             for (int d = 1; d <= exponent; ++d) {
               int offset = pow(2, d - 1);
               kernScan << < fullBlocksPerGrid, blockSize >> > (paddedLength, offset, dev_odata, dev_idata);
@@ -103,7 +79,6 @@ namespace StreamCompaction {
 
             cudaFree(dev_idata);
             cudaFree(dev_odata);
-            cudaFree(dev_unpadded_idata);
         }
 
         // oTrueData[index] is 1 if idata[index] has bit 1 at position i

From b9db98105eeb8c6307d8bf21a1194f8ec86bcb09 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 22:26:04 -0400
Subject: [PATCH 13/17] update readme

---
 README.md                                     | 118 +++++++++++++++++-
 ...plementations on Different Array Sizes.svg |   1 +
 src/main.cpp                                  |   2 +-
 3 files changed, 118 insertions(+), 3 deletions(-)
 create mode 100644 img/Performance of Scan Implementations on Different Array Sizes.svg

diff --git a/README.md b/README.md
index a4aff72..bcf05a1 100644
--- a/README.md
+++ b/README.md
@@ -12,6 +12,120 @@ Tested on AORUS 15P XD laptop with specs:
 - NVIDIA GeForce RTX 3070 Laptop GPU  
 
 I implemented the following parallel algorithms on the GPU and benchmarked them against my own CPU implementations and Thrust on the GPU:  
-- Scan
+- Naive Scan
+- Work efficient scan
 - Stream compaction
-- Radix sort
\ No newline at end of file
+- Radix sort
+
+I roughly optimized the block size for each algorithm by seeing what block size performed the fastest on arrays of size 2^22 (approx 4 million).
+
+| Block Size |	Runtime - Naive (ms)	| Runtime - Work efficient (ms) |
+| ----------- | ----------- | ----------- |
+32 |	4.21734 |	1.32106
+64	|2.11395	|1.36259
+128	|2.09267|	1.26221
+256|	2.09258	|1.28563
+384	|2.11395	|1.3327
+768	|2.11405|	1.26701
+
+Performance was pretty similar for most block sizes, but started to suffer for both naive and work efficient at around 32 or 64 threads per block. In this case, I decided to use a block size of 128 threads to compare the algorithms on different array sizes.  
+
+![](img/Performance%20of%20Scan%20Implementations%20on%20Different%20Array%20Sizes.svg)
+
+| Array size | CPU (m/s)    | Naive (m/s)  | Work efficient (m/s) | Thrust (m/s)  |
+|------------|---------|----------|----------------|----------|
+| 65536      | 0.1023  | 0.158464 | 0.266592       | 0.045472 |
+| 262144     | 0.399   | 0.2616   | 0.33888        | 0.194144 |
+| 1048576    | 1.6835  | 0.636288 | 0.472416       | 0.351648 |
+| 4194304    | 6.392   | 2.20544  | 1.27302        | 0.523776 |
+| 16777216   | 25.5751 | 8.98938  | 4.05302        | 1.09213  |
+| 67108864   | 100.736 | 38.8708  | 15.4414        | 2.14362  |
+| 268435456  | 410.365 | 169.486  | 60.6265        | 6.23341  |
+
+### Extra credit
+
+#### Performance
+My work efficient scan halves the total number of threads launched each iteration, so the work efficient scan is faster at than naive and CPU implementation at array sizes of 2^18 and larger. This means less threads are idling and taking up space on the GPU multi-processors.
+
+#### Radix sort
+
+I implemented Radix sort on the CPU, wrote two test cases, and added Radix sort on the GPU, which calls my work efficient scan. The functions to look at are:
+- `radixSort` in `naive.cu`
+- `radixSort` in `cpu.cu`
+
+A few notes: you can pass in the number of bits you want to sort by, which should be `ilog2ceil(MAX_ARRAY_ELEMENT_VALUE)`. Also, I assumed for simplicity that each element is a positive integer (although still using int and not unsigned int types) so I can just use a bitmask to compact the arrays. Finally, to test, the array size should not be too close to 2^31 because of integer overflow issues...  
+
+#### Sample program output
+Settings: `blockSize` = 128, `SIZE` = 1 << 20
+
+```
+
+****************
+** SCAN TESTS **
+****************
+    [   8  40  25  15  22   4  39  46  44  28   2  44  33 ...  30   0 ]
+==== cpu scan, power-of-two ====
+   elapsed time: 1.5927ms    (std::chrono Measured)
+    [   0   8  48  73  88 110 114 153 199 243 271 273 317 ... 25674057 25674087 ]
+==== cpu scan, non-power-of-two ====
+   elapsed time: 1.7177ms    (std::chrono Measured)
+    [   0   8  48  73  88 110 114 153 199 243 271 273 317 ... 25673959 25673964 ]
+    passed
+==== naive scan, power-of-two ====
+   elapsed time: 0.634144ms    (CUDA Measured)
+    passed
+==== naive scan, non-power-of-two ====
+   elapsed time: 0.602528ms    (CUDA Measured)
+    passed
+==== work-efficient scan, power-of-two ====
+   elapsed time: 0.457504ms    (CUDA Measured)
+    passed
+==== work-efficient scan, non-power-of-two ====
+   elapsed time: 0.479328ms    (CUDA Measured)
+    passed
+==== thrust scan, power-of-two ====
+   elapsed time: 0.340864ms    (CUDA Measured)
+    passed
+==== thrust scan, non-power-of-two ====
+   elapsed time: 0.279648ms    (CUDA Measured)
+    passed
+
+*****************************
+** STREAM COMPACTION TESTS **
+*****************************
+    [   2   0   2   2   3   1   2   1   1   2   2   0   3 ...   0   0 ]
+==== cpu compact without scan, power-of-two ====
+   elapsed time: 2.2561ms    (std::chrono Measured)
+    [   2   2   2   3   1   2   1   1   2   2   3   1   2 ...   1   1 ]
+    passed
+==== cpu compact without scan, non-power-of-two ====
+   elapsed time: 2.1692ms    (std::chrono Measured)
+    [   2   2   2   3   1   2   1   1   2   2   3   1   2 ...   1   1 ]
+    passed
+==== cpu compact with scan ====
+   elapsed time: 5.498ms    (std::chrono Measured)
+    [   2   2   2   3   1   2   1   1   2   2   3   1   2 ...   1   1 ]
+    passed
+==== work-efficient compact, power-of-two ====
+   elapsed time: 0.541312ms    (CUDA Measured)
+    passed
+==== work-efficient compact, non-power-of-two ====
+   elapsed time: 0.490016ms    (CUDA Measured)
+    passed
+
+*****************************
+** RADIX SORT TESTS **
+*****************************
+    [ 23914 19768 8186 18274 8083 25457 19314 28109 22549 17382 25706 27432 8575 ... 4256   0 ]
+==== cpu radix sort, power-of-two ====
+    [   0   0   0   0   0   0   0   0   0   0   0   0   0 ... 32767 32767 ]
+==== radix sort, power of two ====
+   elapsed time: 0.490016ms    (CUDA Measured)
+    passed
+==== cpu radix sort, non-power-of-two ====
+   elapsed time: 109.43ms    (std::chrono Measured)
+    [   0   0   0   0   0   0   0   0   0   0   0   0   0 ... 32767 32767 ]
+==== radix sort, non-power of two ====
+   elapsed time: 0.490016ms    (CUDA Measured)
+    passed
+```
\ No newline at end of file
diff --git a/img/Performance of Scan Implementations on Different Array Sizes.svg b/img/Performance of Scan Implementations on Different Array Sizes.svg
new file mode 100644
index 0000000..db0a1e5
--- /dev/null
+++ b/img/Performance of Scan Implementations on Different Array Sizes.svg	
@@ -0,0 +1 @@
+<svg version="1.1" viewBox="0.0 0.0 682.0 467.0" fill="none" stroke="none" stroke-linecap="square" stroke-miterlimit="10" width="682" height="467" xmlns:xlink="http://www.w3.org/1999/xlink" xmlns="http://www.w3.org/2000/svg"><path fill="#ffffff" d="M0 0L682.0 0L682.0 467.0L0 467.0L0 0Z" fill-rule="nonzero"/><path stroke="#333333" stroke-width="1.0" stroke-linecap="butt" d="M89.5 69.5L89.5 391.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M171.5 69.5L171.5 391.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M254.5 69.5L254.5 391.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M337.5 69.5L337.5 391.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M420.5 69.5L420.5 391.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M502.5 69.5L502.5 391.5" fill-rule="nonzero"/><path stroke="#333333" stroke-width="1.0" stroke-linecap="butt" d="M82.5 391.5L540.5 391.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M82.5 327.5L540.5 327.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M82.5 262.5L540.5 262.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M82.5 198.5L540.5 198.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M82.5 134.5L540.5 134.5" fill-rule="nonzero"/><path stroke="#cccccc" stroke-width="1.0" stroke-linecap="butt" d="M82.5 69.5L540.5 69.5" fill-rule="nonzero"/><clipPath id="id_0"><path d="M82.35 69.916664L540.26776 69.916664L540.26776 391.65L82.35 391.65L82.35 69.916664Z" clip-rule="nonzero"/></clipPath><path stroke="#4285f4" stroke-width="2.0" stroke-linecap="butt" clip-path="url(#id_0)" d="M89.142334 391.58417L89.46786 391.39325L90.76993 390.5667L95.97823 387.53696L116.81143 375.19327L200.14423 326.82974L533.47546 127.5938" fill-rule="nonzero"/><path stroke="#ea4335" stroke-width="2.0" stroke-linecap="butt" clip-path="url(#id_0)" d="M89.142334 391.54803L89.46786 391.48166L90.76993 391.24057L95.97823 390.23087L116.81143 385.86563L200.14423 366.63794L533.47546 282.5914" fill-rule="nonzero"/><path stroke="#fbbc04" stroke-width="2.0" stroke-linecap="butt" clip-path="url(#id_0)" d="M89.142334 391.47845L89.46786 391.43195L90.76993 391.346L95.97823 390.83084L116.81143 389.04202L200.14423 381.71396L533.47546 352.63885" fill-rule="nonzero"/><path stroke="#34a853" stroke-width="2.0" stroke-linecap="butt" clip-path="url(#id_0)" d="M89.142334 391.62073L89.46786 391.5251L90.76993 391.42374L95.97823 391.31296L116.81143 390.94727L200.14423 390.27066L533.47546 387.639" fill-rule="nonzero"/><path fill="#4285f4" clip-path="url(#id_0)" d="M93.642334 391.58417C93.642334 394.06946 91.62762 396.08417 89.142334 396.08417C86.65706 396.08417 84.642334 394.06946 84.642334 391.58417C84.642334 389.09888 86.65706 387.08417 89.142334 387.08417C91.62762 387.08417 93.642334 389.09888 93.642334 391.58417Z" fill-rule="nonzero"/><path fill="#4285f4" clip-path="url(#id_0)" d="M93.96786 391.39325C93.96786 393.87854 91.95314 395.89325 89.46786 395.89325C86.982574 395.89325 84.96786 393.87854 84.96786 391.39325C84.96786 388.908 86.982574 386.89325 89.46786 386.89325C91.95314 386.89325 93.96786 388.908 93.96786 391.39325Z" fill-rule="nonzero"/><path fill="#4285f4" clip-path="url(#id_0)" d="M95.26993 390.5667C95.26993 393.052 93.25521 395.0667 90.76993 395.0667C88.28465 395.0667 86.26993 393.052 86.26993 390.5667C86.26993 388.08145 88.28465 386.0667 90.76993 386.0667C93.25521 386.0667 95.26993 388.08145 95.26993 390.5667Z" fill-rule="nonzero"/><path fill="#4285f4" clip-path="url(#id_0)" d="M100.47823 387.53696C100.47823 390.02225 98.46351 392.03696 95.97823 392.03696C93.49295 392.03696 91.47823 390.02225 91.47823 387.53696C91.47823 385.05167 93.49295 383.03696 95.97823 383.03696C98.46351 383.03696 100.47823 385.05167 100.47823 387.53696Z" fill-rule="nonzero"/><path fill="#4285f4" clip-path="url(#id_0)" d="M121.31143 375.19327C121.31143 377.67856 119.296715 379.69327 116.81143 379.69327C114.32615 379.69327 112.31143 377.67856 112.31143 375.19327C112.31143 372.708 114.32615 370.69327 116.81143 370.69327C119.296715 370.69327 121.31143 372.708 121.31143 375.19327Z" fill-rule="nonzero"/><path fill="#4285f4" clip-path="url(#id_0)" d="M204.64423 326.82974C204.64423 329.31503 202.62952 331.32974 200.14423 331.32974C197.65895 331.32974 195.64423 329.31503 195.64423 326.82974C195.64423 324.34445 197.65895 322.32974 200.14423 322.32974C202.62952 322.32974 204.64423 324.34445 204.64423 326.82974Z" fill-rule="nonzero"/><path fill="#4285f4" clip-path="url(#id_0)" d="M537.97546 127.5938C537.97546 130.07909 535.9607 132.0938 533.47546 132.0938C530.9902 132.0938 528.97546 130.07909 528.97546 127.5938C528.97546 125.10852 530.9902 123.0938 533.47546 123.0938C535.9607 123.0938 537.97546 125.10852 537.97546 127.5938Z" fill-rule="nonzero"/><path fill="#ea4335" clip-path="url(#id_0)" d="M93.642334 391.54803C93.642334 394.03333 91.62762 396.04803 89.142334 396.04803C86.65706 396.04803 84.642334 394.03333 84.642334 391.54803C84.642334 389.06274 86.65706 387.04803 89.142334 387.04803C91.62762 387.04803 93.642334 389.06274 93.642334 391.54803Z" fill-rule="nonzero"/><path fill="#ea4335" clip-path="url(#id_0)" d="M93.96786 391.48166C93.96786 393.96695 91.95314 395.98166 89.46786 395.98166C86.982574 395.98166 84.96786 393.96695 84.96786 391.48166C84.96786 388.9964 86.982574 386.98166 89.46786 386.98166C91.95314 386.98166 93.96786 388.9964 93.96786 391.48166Z" fill-rule="nonzero"/><path fill="#ea4335" clip-path="url(#id_0)" d="M95.26993 391.24057C95.26993 393.72586 93.25521 395.74057 90.76993 395.74057C88.28465 395.74057 86.26993 393.72586 86.26993 391.24057C86.26993 388.75528 88.28465 386.74057 90.76993 386.74057C93.25521 386.74057 95.26993 388.75528 95.26993 391.24057Z" fill-rule="nonzero"/><path fill="#ea4335" clip-path="url(#id_0)" d="M100.47823 390.23087C100.47823 392.71616 98.46351 394.73087 95.97823 394.73087C93.49295 394.73087 91.47823 392.71616 91.47823 390.23087C91.47823 387.7456 93.49295 385.73087 95.97823 385.73087C98.46351 385.73087 100.47823 387.7456 100.47823 390.23087Z" fill-rule="nonzero"/><path fill="#ea4335" clip-path="url(#id_0)" d="M121.31143 385.86563C121.31143 388.35092 119.296715 390.36563 116.81143 390.36563C114.32615 390.36563 112.31143 388.35092 112.31143 385.86563C112.31143 383.38034 114.32615 381.36563 116.81143 381.36563C119.296715 381.36563 121.31143 383.38034 121.31143 385.86563Z" fill-rule="nonzero"/><path fill="#ea4335" clip-path="url(#id_0)" d="M204.64423 366.63794C204.64423 369.12323 202.62952 371.13794 200.14423 371.13794C197.65895 371.13794 195.64423 369.12323 195.64423 366.63794C195.64423 364.15265 197.65895 362.13794 200.14423 362.13794C202.62952 362.13794 204.64423 364.15265 204.64423 366.63794Z" fill-rule="nonzero"/><path fill="#ea4335" clip-path="url(#id_0)" d="M537.97546 282.5914C537.97546 285.0767 535.9607 287.0914 533.47546 287.0914C530.9902 287.0914 528.97546 285.0767 528.97546 282.5914C528.97546 280.10614 530.9902 278.0914 533.47546 278.0914C535.9607 278.0914 537.97546 280.10614 537.97546 282.5914Z" fill-rule="nonzero"/><path fill="#fbbc04" clip-path="url(#id_0)" d="M93.642334 391.47845C93.642334 393.96375 91.62762 395.97845 89.142334 395.97845C86.65706 395.97845 84.642334 393.96375 84.642334 391.47845C84.642334 388.99316 86.65706 386.97845 89.142334 386.97845C91.62762 386.97845 93.642334 388.99316 93.642334 391.47845Z" fill-rule="nonzero"/><path fill="#fbbc04" clip-path="url(#id_0)" d="M93.96786 391.43195C93.96786 393.91724 91.95314 395.93195 89.46786 395.93195C86.982574 395.93195 84.96786 393.91724 84.96786 391.43195C84.96786 388.94666 86.982574 386.93195 89.46786 386.93195C91.95314 386.93195 93.96786 388.94666 93.96786 391.43195Z" fill-rule="nonzero"/><path fill="#fbbc04" clip-path="url(#id_0)" d="M95.26993 391.346C95.26993 393.8313 93.25521 395.846 90.76993 395.846C88.28465 395.846 86.26993 393.8313 86.26993 391.346C86.26993 388.86075 88.28465 386.846 90.76993 386.846C93.25521 386.846 95.26993 388.86075 95.26993 391.346Z" fill-rule="nonzero"/><path fill="#fbbc04" clip-path="url(#id_0)" d="M100.47823 390.83084C100.47823 393.31613 98.46351 395.33084 95.97823 395.33084C93.49295 395.33084 91.47823 393.31613 91.47823 390.83084C91.47823 388.34558 93.49295 386.33084 95.97823 386.33084C98.46351 386.33084 100.47823 388.34558 100.47823 390.83084Z" fill-rule="nonzero"/><path fill="#fbbc04" clip-path="url(#id_0)" d="M121.31143 389.04202C121.31143 391.5273 119.296715 393.54202 116.81143 393.54202C114.32615 393.54202 112.31143 391.5273 112.31143 389.04202C112.31143 386.55673 114.32615 384.54202 116.81143 384.54202C119.296715 384.54202 121.31143 386.55673 121.31143 389.04202Z" fill-rule="nonzero"/><path fill="#fbbc04" clip-path="url(#id_0)" d="M204.64423 381.71396C204.64423 384.19925 202.62952 386.21396 200.14423 386.21396C197.65895 386.21396 195.64423 384.19925 195.64423 381.71396C195.64423 379.2287 197.65895 377.21396 200.14423 377.21396C202.62952 377.21396 204.64423 379.2287 204.64423 381.71396Z" fill-rule="nonzero"/><path fill="#fbbc04" clip-path="url(#id_0)" d="M537.97546 352.63885C537.97546 355.12415 535.9607 357.13885 533.47546 357.13885C530.9902 357.13885 528.97546 355.12415 528.97546 352.63885C528.97546 350.1536 530.9902 348.13885 533.47546 348.13885C535.9607 348.13885 537.97546 350.1536 537.97546 352.63885Z" fill-rule="nonzero"/><path fill="#34a853" clip-path="url(#id_0)" d="M93.642334 391.62073C93.642334 394.10602 91.62762 396.12073 89.142334 396.12073C86.65706 396.12073 84.642334 394.10602 84.642334 391.62073C84.642334 389.13547 86.65706 387.12073 89.142334 387.12073C91.62762 387.12073 93.642334 389.13547 93.642334 391.62073Z" fill-rule="nonzero"/><path fill="#34a853" clip-path="url(#id_0)" d="M93.96786 391.5251C93.96786 394.01035 91.95314 396.0251 89.46786 396.0251C86.982574 396.0251 84.96786 394.01035 84.96786 391.5251C84.96786 389.0398 86.982574 387.0251 89.46786 387.0251C91.95314 387.0251 93.96786 389.0398 93.96786 391.5251Z" fill-rule="nonzero"/><path fill="#34a853" clip-path="url(#id_0)" d="M95.26993 391.42374C95.26993 393.909 93.25521 395.92374 90.76993 395.92374C88.28465 395.92374 86.26993 393.909 86.26993 391.42374C86.26993 388.93845 88.28465 386.92374 90.76993 386.92374C93.25521 386.92374 95.26993 388.93845 95.26993 391.42374Z" fill-rule="nonzero"/><path fill="#34a853" clip-path="url(#id_0)" d="M100.47823 391.31296C100.47823 393.79825 98.46351 395.81296 95.97823 395.81296C93.49295 395.81296 91.47823 393.79825 91.47823 391.31296C91.47823 388.8277 93.49295 386.81296 95.97823 386.81296C98.46351 386.81296 100.47823 388.8277 100.47823 391.31296Z" fill-rule="nonzero"/><path fill="#34a853" clip-path="url(#id_0)" d="M121.31143 390.94727C121.31143 393.43253 119.296715 395.44727 116.81143 395.44727C114.32615 395.44727 112.31143 393.43253 112.31143 390.94727C112.31143 388.46198 114.32615 386.44727 116.81143 386.44727C119.296715 386.44727 121.31143 388.46198 121.31143 390.94727Z" fill-rule="nonzero"/><path fill="#34a853" clip-path="url(#id_0)" d="M204.64423 390.27066C204.64423 392.75592 202.62952 394.77066 200.14423 394.77066C197.65895 394.77066 195.64423 392.75592 195.64423 390.27066C195.64423 387.78537 197.65895 385.77066 200.14423 385.77066C202.62952 385.77066 204.64423 387.78537 204.64423 390.27066Z" fill-rule="nonzero"/><path fill="#34a853" clip-path="url(#id_0)" d="M537.97546 387.639C537.97546 390.1243 535.9607 392.139 533.47546 392.139C530.9902 392.139 528.97546 390.1243 528.97546 387.639C528.97546 385.15372 530.9902 383.139 533.47546 383.139C535.9607 383.139 537.97546 385.15372 537.97546 387.639Z" fill-rule="nonzero"/><path fill="#000000" d="M285.29327 440.65L288.59015 432.05624L289.8089 432.05624L293.32452 440.65L292.04327 440.65L291.02765 438.04062L287.44952 438.04062L286.4964 440.65L285.29327 440.65ZM287.77765 437.11874L290.6839 437.11874L289.79327 434.74374Q289.3714 433.66562 289.1839 432.9625Q289.01202 433.79062 288.71515 434.60312L287.77765 437.11874ZM293.09015 440.65L293.09015 434.43124L294.04327 434.43124L294.04327 435.36874Q294.40265 434.7125 294.69952 434.50937Q295.01202 434.29062 295.38702 434.29062Q295.91827 434.29062 296.46515 434.61874L296.10577 435.60312Q295.71515 435.36874 295.34015 435.36874Q294.98077 435.36874 294.69952 435.5875Q294.4339 435.79062 294.32452 436.16562Q294.13702 436.72812 294.13702 437.38437L294.13702 440.65L293.09015 440.65ZM297.09015 440.65L297.09015 434.43124L298.04327 434.43124L298.04327 435.36874Q298.40265 434.7125 298.69952 434.50937Q299.01202 434.29062 299.38702 434.29062Q299.91827 434.29062 300.46515 434.61874L300.10577 435.60312Q299.71515 435.36874 299.34015 435.36874Q298.98077 435.36874 298.69952 435.5875Q298.4339 435.79062 298.32452 436.16562Q298.13702 436.72812 298.13702 437.38437L298.13702 440.65L297.09015 440.65ZM305.16827 439.88437Q304.57452 440.38437 304.02765 440.5875Q303.4964 440.79062 302.8714 440.79062Q301.84015 440.79062 301.29327 440.29062Q300.7464 439.79062 300.7464 439.00937Q300.7464 438.55624 300.94952 438.18124Q301.15265 437.79062 301.4964 437.57187Q301.84015 437.3375 302.26202 437.22812Q302.5589 437.13437 303.19952 437.05624Q304.46515 436.91562 305.07452 436.69687Q305.07452 436.47812 305.07452 436.43124Q305.07452 435.775 304.77765 435.50937Q304.3714 435.16562 303.57452 435.16562Q302.84015 435.16562 302.48077 435.43124Q302.1214 435.68124 301.94952 436.3375L300.91827 436.19687Q301.0589 435.54062 301.38702 435.13437Q301.71515 434.72812 302.32452 434.50937Q302.9339 434.29062 303.73077 434.29062Q304.52765 434.29062 305.02765 434.47812Q305.52765 434.66562 305.76202 434.94687Q305.9964 435.22812 306.09015 435.66562Q306.13702 435.93124 306.13702 436.63437L306.13702 438.04062Q306.13702 439.50937 306.19952 439.9Q306.27765 440.29062 306.48077 440.65L305.3714 440.65Q305.21515 440.32187 305.16827 439.88437ZM305.07452 437.525Q304.4964 437.75937 303.35577 437.93124Q302.69952 438.025 302.4339 438.15Q302.16827 438.25937 302.01202 438.47812Q301.8714 438.69687 301.8714 438.97812Q301.8714 439.4 302.1839 439.68124Q302.51202 439.9625 303.1214 439.9625Q303.73077 439.9625 304.19952 439.69687Q304.6839 439.43124 304.90265 438.9625Q305.07452 438.60312 305.07452 437.91562L305.07452 437.525ZM308.0589 443.04062L307.9339 442.05624Q308.27765 442.15 308.54327 442.15Q308.88702 442.15 309.09015 442.025Q309.3089 441.91562 309.44952 441.7125Q309.54327 441.54062 309.77765 440.91562Q309.79327 440.8375 309.8714 440.66562L307.4964 434.43124L308.63702 434.43124L309.9339 438.025Q310.1839 438.7125 310.38702 439.47812Q310.57452 438.74374 310.82452 438.05624L312.15265 434.43124L313.19952 434.43124L310.84015 440.75937Q310.44952 441.775 310.2464 442.16562Q309.96515 442.69687 309.59015 442.93124Q309.23077 443.18124 308.73077 443.18124Q308.4339 443.18124 308.0589 443.04062ZM315.6839 438.79062L316.71515 438.63437Q316.8089 439.25937 317.19952 439.5875Q317.60577 439.91562 318.34015 439.91562Q319.0589 439.91562 319.40265 439.63437Q319.76202 439.3375 319.76202 438.93124Q319.76202 438.57187 319.44952 438.36874Q319.23077 438.22812 318.3714 438.00937Q317.21515 437.7125 316.76202 437.50937Q316.32452 437.29062 316.09015 436.91562Q315.85577 436.54062 315.85577 436.07187Q315.85577 435.66562 316.04327 435.30624Q316.23077 434.94687 316.5589 434.7125Q316.8089 434.54062 317.23077 434.41562Q317.65265 434.29062 318.15265 434.29062Q318.8714 434.29062 319.41827 434.50937Q319.98077 434.7125 320.2464 435.07187Q320.51202 435.43124 320.60577 436.025L319.57452 436.16562Q319.51202 435.69687 319.16827 435.43124Q318.84015 435.15 318.21515 435.15Q317.4964 435.15 317.1839 435.4Q316.8714 435.63437 316.8714 435.9625Q316.8714 436.16562 316.9964 436.32187Q317.13702 436.49374 317.40265 436.60312Q317.5589 436.66562 318.34015 436.86874Q319.46515 437.18124 319.90265 437.36874Q320.34015 437.55624 320.59015 437.91562Q320.84015 438.275 320.84015 438.82187Q320.84015 439.35312 320.52765 439.82187Q320.23077 440.275 319.65265 440.54062Q319.07452 440.79062 318.34015 440.79062Q317.1214 440.79062 316.48077 440.29062Q315.85577 439.775 315.6839 438.79062ZM323.10577 433.275L323.10577 432.05624L324.16827 432.05624L324.16827 433.275L323.10577 433.275ZM323.10577 440.65L323.10577 434.43124L324.16827 434.43124L324.16827 440.65L323.10577 440.65ZM325.54327 440.65L325.54327 439.79062L329.51202 435.24374Q328.82452 435.275 328.3089 435.275L325.77765 435.275L325.77765 434.43124L330.8714 434.43124L330.8714 435.11874L327.4964 439.07187L326.84015 439.79062Q327.5589 439.74374 328.16827 439.74374L331.0589 439.74374L331.0589 440.65L325.54327 440.65ZM335.35577 438.65L336.44952 438.775Q336.19952 439.72812 335.4964 440.25937Q334.79327 440.79062 333.71515 440.79062Q332.35577 440.79062 331.54327 439.94687Q330.7464 439.10312 330.7464 437.5875Q330.7464 436.025 331.5589 435.16562Q332.3714 434.29062 333.65265 434.29062Q334.90265 434.29062 335.6839 435.13437Q336.48077 435.97812 336.48077 437.525Q336.48077 437.61874 336.48077 437.80624L331.84015 437.80624Q331.90265 438.8375 332.41827 439.38437Q332.9339 439.91562 333.71515 439.91562Q334.29327 439.91562 334.69952 439.61874Q335.1214 439.30624 335.35577 438.65ZM331.90265 436.94687L335.3714 436.94687Q335.3089 436.15 334.98077 435.75937Q334.46515 435.15 333.66827 435.15Q332.9339 435.15 332.4339 435.63437Q331.94952 436.11874 331.90265 436.94687Z" fill-rule="nonzero"/><path fill="#000000" d="M34.349995 275.84583L25.756245 275.84583L25.756245 272.03333Q25.756245 270.87708 25.99062 270.28333Q26.224995 269.68958 26.818745 269.3302Q27.39687 268.97083 28.099995 268.97083Q29.02187 268.97083 29.64687 269.56458Q30.27187 270.15833 30.443745 271.3927Q30.662495 270.93958 30.86562 270.7052Q31.318745 270.2052 32.006245 269.7677L34.349995 268.2677L34.349995 269.7052L32.568745 270.8302Q31.787495 271.3302 31.381245 271.65833Q30.974995 271.97083 30.818745 272.23645Q30.64687 272.48645 30.568745 272.75208Q30.537495 272.93958 30.537495 273.37708L30.537495 274.7052L34.349995 274.7052L34.349995 275.84583ZM29.55312 274.7052L29.55312 272.25208Q29.55312 271.48645 29.39687 271.04895Q29.224995 270.59583 28.881245 270.37708Q28.52187 270.1427 28.099995 270.1427Q27.49062 270.1427 27.099995 270.59583Q26.70937 271.03333 26.70937 271.98645L26.70937 274.7052L29.55312 274.7052ZM34.349995 262.90833L33.42812 262.90833Q34.49062 263.6427 34.49062 264.8927Q34.49062 265.43958 34.287495 265.92395Q34.068745 266.3927 33.756245 266.62708Q33.42812 266.86145 32.95937 266.9552Q32.662495 267.0177 31.974995 267.0177L28.131245 267.0177L28.131245 265.9552L31.58437 265.9552Q32.39687 265.9552 32.693745 265.8927Q33.099995 265.79895 33.349995 265.48645Q33.58437 265.15833 33.58437 264.67395Q33.58437 264.2052 33.349995 263.79895Q33.099995 263.37708 32.67812 263.2052Q32.256245 263.0177 31.45937 263.0177L28.131245 263.0177L28.131245 261.97083L34.349995 261.97083L34.349995 262.90833ZM34.349995 259.98645L28.131245 259.98645L28.131245 259.04895L29.006245 259.04895Q27.99062 258.36145 27.99062 257.06458Q27.99062 256.50208 28.193745 256.03333Q28.39687 255.54897 28.724995 255.31459Q29.05312 255.08022 29.49062 254.98647Q29.787495 254.93959 30.52187 254.93959L34.349995 254.93959L34.349995 255.98647L30.568745 255.98647Q29.912495 255.98647 29.599995 256.11145Q29.287495 256.23645 29.099995 256.54895Q28.89687 256.86145 28.89687 257.28333Q28.89687 257.9552 29.33437 258.4552Q29.756245 258.93958 30.943745 258.93958L34.349995 258.93958L34.349995 259.98645ZM33.412495 250.68959L34.33437 250.53334Q34.42812 250.98647 34.42812 251.33022Q34.42812 251.90834 34.256245 252.22084Q34.068745 252.53334 33.77187 252.67397Q33.474995 252.79897 32.52187 252.79897L28.943745 252.79897L28.943745 253.56459L28.131245 253.56459L28.131245 252.79897L26.58437 252.79897L25.95937 251.75209L28.131245 251.75209L28.131245 250.68959L28.943745 250.68959L28.943745 251.75209L32.58437 251.75209Q33.037495 251.75209 33.162495 251.70522Q33.287495 251.64272 33.36562 251.51772Q33.443745 251.39272 33.443745 251.15834Q33.443745 250.97084 33.412495 250.68959ZM26.974995 249.98647L25.756245 249.98647L25.756245 248.92397L26.974995 248.92397L26.974995 249.98647ZM34.349995 249.98647L28.131245 249.98647L28.131245 248.92397L34.349995 248.92397L34.349995 249.98647ZM34.349995 246.98647L28.131245 246.98647L28.131245 246.04897L29.006245 246.04897Q28.537495 245.75209 28.27187 245.26772Q27.99062 244.78334 27.99062 244.15834Q27.99062 243.47084 28.27187 243.03334Q28.55312 242.58022 29.068745 242.40834Q27.99062 241.65834 27.99062 240.48647Q27.99062 239.54897 28.506245 239.06459Q29.006245 238.56459 30.08437 238.56459L34.349995 238.56459L34.349995 239.61147L30.42812 239.61147Q29.80312 239.61147 29.52187 239.72084Q29.24062 239.81459 29.068745 240.08022Q28.89687 240.34584 28.89687 240.72084Q28.89687 241.37709 29.33437 241.81459Q29.77187 242.23647 30.74062 242.23647L34.349995 242.23647L34.349995 243.29897L30.30312 243.29897Q29.599995 243.29897 29.256245 243.56459Q28.89687 243.81459 28.89687 244.39272Q28.89687 244.84584 29.131245 245.22084Q29.36562 245.59584 29.818745 245.76772Q30.27187 245.93959 31.11562 245.93959L34.349995 245.93959L34.349995 246.98647ZM32.349995 231.73647L32.474995 230.64272Q33.42812 230.89272 33.95937 231.59584Q34.49062 232.29897 34.49062 233.37709Q34.49062 234.73647 33.64687 235.54897Q32.80312 236.34584 31.287495 236.34584Q29.724995 236.34584 28.86562 235.53334Q27.99062 234.72084 27.99062 233.43959Q27.99062 232.18959 28.83437 231.40834Q29.67812 230.61147 31.224995 230.61147Q31.318745 230.61147 31.506245 230.61147L31.506245 235.25209Q32.537495 235.18959 33.08437 234.67397Q33.61562 234.15834 33.61562 233.37709Q33.61562 232.79897 33.318745 232.39272Q33.006245 231.97084 32.349995 231.73647ZM30.64687 235.18959L30.64687 231.72084Q29.849995 231.78334 29.45937 232.11147Q28.849995 232.62709 28.849995 233.42397Q28.849995 234.15834 29.33437 234.65834Q29.818745 235.14272 30.64687 235.18959ZM36.881245 223.97084Q35.77187 224.84584 34.30312 225.45522Q32.818745 226.04897 31.24062 226.04897Q29.849995 226.04897 28.568745 225.61147Q27.08437 225.08022 25.61562 223.97084L25.61562 223.22084Q26.83437 223.92397 27.349995 224.15834Q28.162495 224.51772 29.037495 224.72084Q30.131245 224.97084 31.24062 224.97084Q34.068745 224.97084 36.881245 223.22084L36.881245 223.97084ZM34.349995 221.98647L28.131245 221.98647L28.131245 221.04897L29.006245 221.04897Q28.537495 220.75209 28.27187 220.26772Q27.99062 219.78334 27.99062 219.15834Q27.99062 218.47084 28.27187 218.03334Q28.55312 217.58022 29.068745 217.40834Q27.99062 216.65834 27.99062 215.48647Q27.99062 214.54897 28.506245 214.06459Q29.006245 213.56459 30.08437 213.56459L34.349995 213.56459L34.349995 214.61147L30.42812 214.61147Q29.80312 214.61147 29.52187 214.72084Q29.24062 214.81459 29.068745 215.08022Q28.89687 215.34584 28.89687 215.72084Q28.89687 216.37709 29.33437 216.81459Q29.77187 217.23647 30.74062 217.23647L34.349995 217.23647L34.349995 218.29897L30.30312 218.29897Q29.599995 218.29897 29.256245 218.56459Q28.89687 218.81459 28.89687 219.39272Q28.89687 219.84584 29.131245 220.22084Q29.36562 220.59584 29.818745 220.76772Q30.27187 220.93959 31.11562 220.93959L34.349995 220.93959L34.349995 221.98647ZM34.49062 211.78334L25.61562 209.29897L25.61562 208.45522L34.49062 210.93959L34.49062 211.78334ZM32.49062 208.40834L32.33437 207.37709Q32.95937 207.28334 33.287495 206.89272Q33.61562 206.48647 33.61562 205.75209Q33.61562 205.03334 33.33437 204.68959Q33.037495 204.33022 32.631245 204.33022Q32.27187 204.33022 32.068745 204.64272Q31.92812 204.86147 31.70937 205.72084Q31.412495 206.87709 31.20937 207.33022Q30.99062 207.76772 30.61562 208.00209Q30.24062 208.23647 29.77187 208.23647Q29.36562 208.23647 29.006245 208.04897Q28.64687 207.86147 28.412495 207.53334Q28.24062 207.28334 28.11562 206.86147Q27.99062 206.43959 27.99062 205.93959Q27.99062 205.22084 28.20937 204.67397Q28.412495 204.11147 28.77187 203.84584Q29.131245 203.58022 29.724995 203.48647L29.86562 204.51772Q29.39687 204.58022 29.131245 204.92397Q28.849995 205.25209 28.849995 205.87709Q28.849995 206.59584 29.099995 206.90834Q29.33437 207.22084 29.662495 207.22084Q29.86562 207.22084 30.02187 207.09584Q30.193745 206.95522 30.30312 206.68959Q30.36562 206.53334 30.568745 205.75209Q30.881245 204.62709 31.068745 204.18959Q31.256245 203.75209 31.61562 203.50209Q31.974995 203.25209 32.52187 203.25209Q33.05312 203.25209 33.52187 203.56459Q33.974995 203.86147 34.24062 204.43959Q34.49062 205.01772 34.49062 205.75209Q34.49062 206.97084 33.99062 207.61147Q33.474995 208.23647 32.49062 208.40834ZM36.881245 200.29897L36.881245 201.04897Q34.068745 199.29897 31.24062 199.29897Q30.14687 199.29897 29.05312 199.54897Q28.17812 199.75209 27.36562 200.11147Q26.849995 200.34584 25.61562 201.04897L25.61562 200.29897Q27.08437 199.20522 28.568745 198.67397Q29.849995 198.22084 31.24062 198.22084Q32.818745 198.22084 34.30312 198.83022Q35.77187 199.43959 36.881245 200.29897Z" fill-rule="nonzero"/><path fill="#000000" d="M69.85 391.41562Q69.85 389.88437 70.1625 388.9625Q70.475 388.025 71.08437 387.525Q71.70937 387.025 72.64687 387.025Q73.33437 387.025 73.85 387.30624Q74.38125 387.5875 74.725 388.11874Q75.06875 388.63437 75.25625 389.38437Q75.44375 390.13437 75.44375 391.41562Q75.44375 392.93124 75.13125 393.85312Q74.83437 394.775 74.20937 395.29062Q73.6 395.79062 72.64687 395.79062Q71.4125 395.79062 70.69375 394.9Q69.85 393.8375 69.85 391.41562ZM70.92812 391.41562Q70.92812 393.525 71.42812 394.22812Q71.92812 394.93124 72.64687 394.93124Q73.38125 394.93124 73.86562 394.22812Q74.36562 393.525 74.36562 391.41562Q74.36562 389.29062 73.86562 388.60312Q73.38125 387.9 72.63125 387.9Q71.9125 387.9 71.475 388.50937Q70.92812 389.29062 70.92812 391.41562Z" fill-rule="nonzero"/><path fill="#000000" d="M59.81875 331.30334L58.771873 331.30334L58.771873 324.5846Q58.38125 324.94397 57.75625 325.31897Q57.146873 325.67834 56.6625 325.85022L56.6625 324.8346Q57.5375 324.41272 58.19375 323.8346Q58.865623 323.24084 59.146873 322.67834L59.81875 322.67834L59.81875 331.30334ZM62.85 327.06897Q62.85 325.53772 63.1625 324.61584Q63.475 323.67834 64.08437 323.17834Q64.70937 322.67834 65.64687 322.67834Q66.33437 322.67834 66.85 322.9596Q67.38125 323.24084 67.725 323.7721Q68.06875 324.28772 68.25625 325.03772Q68.44375 325.78772 68.44375 327.06897Q68.44375 328.5846 68.13125 329.50647Q67.83437 330.42834 67.20937 330.94397Q66.6 331.44397 65.64687 331.44397Q64.4125 331.44397 63.69375 330.55334Q62.85 329.49084 62.85 327.06897ZM63.928123 327.06897Q63.928123 329.17834 64.42812 329.88147Q64.92812 330.5846 65.64687 330.5846Q66.38125 330.5846 66.86562 329.88147Q67.36562 329.17834 67.36562 327.06897Q67.36562 324.94397 66.86562 324.25647Q66.38125 323.55334 65.63125 323.55334Q64.9125 323.55334 64.475 324.16272Q63.928123 324.94397 63.928123 327.06897ZM69.85 327.06897Q69.85 325.53772 70.1625 324.61584Q70.475 323.67834 71.08437 323.17834Q71.70937 322.67834 72.64687 322.67834Q73.33437 322.67834 73.85 322.9596Q74.38125 323.24084 74.725 323.7721Q75.06875 324.28772 75.25625 325.03772Q75.44375 325.78772 75.44375 327.06897Q75.44375 328.5846 75.13125 329.50647Q74.83437 330.42834 74.20937 330.94397Q73.6 331.44397 72.64687 331.44397Q71.4125 331.44397 70.69375 330.55334Q69.85 329.49084 69.85 327.06897ZM70.92812 327.06897Q70.92812 329.17834 71.42812 329.88147Q71.92812 330.5846 72.64687 330.5846Q73.38125 330.5846 73.86562 329.88147Q74.36562 329.17834 74.36562 327.06897Q74.36562 324.94397 73.86562 324.25647Q73.38125 323.55334 72.63125 323.55334Q71.9125 323.55334 71.475 324.16272Q70.92812 324.94397 70.92812 327.06897Z" fill-rule="nonzero"/><path fill="#000000" d="M61.396873 265.94104L61.396873 266.95667L55.709373 266.95667Q55.709373 266.58167 55.834373 266.2223Q56.053123 265.64417 56.521873 265.08167Q57.00625 264.51917 57.9125 263.7848Q59.31875 262.62854 59.803123 261.95667Q60.303123 261.2848 60.303123 260.69104Q60.303123 260.06604 59.85 259.64417Q59.396873 259.20667 58.678123 259.20667Q57.9125 259.20667 57.459373 259.6598Q57.00625 260.1129 56.990623 260.9254L55.9125 260.81604Q56.021873 259.5973 56.740623 258.9723Q57.475 258.33167 58.709373 258.33167Q59.94375 258.33167 60.6625 259.01917Q61.38125 259.70667 61.38125 260.7223Q61.38125 261.2379 61.1625 261.7379Q60.959373 262.2223 60.459373 262.7848Q59.975 263.33167 58.85 264.2848Q57.896873 265.08167 57.615623 265.37854Q57.35 265.6598 57.178123 265.94104L61.396873 265.94104ZM62.85 262.7223Q62.85 261.19104 63.1625 260.26917Q63.475 259.33167 64.08437 258.83167Q64.70937 258.33167 65.64687 258.33167Q66.33437 258.33167 66.85 258.6129Q67.38125 258.89417 67.725 259.4254Q68.06875 259.94104 68.25625 260.69104Q68.44375 261.44104 68.44375 262.7223Q68.44375 264.2379 68.13125 265.1598Q67.83437 266.08167 67.20937 266.5973Q66.6 267.0973 65.64687 267.0973Q64.4125 267.0973 63.69375 266.20667Q62.85 265.14417 62.85 262.7223ZM63.928123 262.7223Q63.928123 264.83167 64.42812 265.5348Q64.92812 266.2379 65.64687 266.2379Q66.38125 266.2379 66.86562 265.5348Q67.36562 264.83167 67.36562 262.7223Q67.36562 260.5973 66.86562 259.9098Q66.38125 259.20667 65.63125 259.20667Q64.9125 259.20667 64.475 259.81604Q63.928123 260.5973 63.928123 262.7223ZM69.85 262.7223Q69.85 261.19104 70.1625 260.26917Q70.475 259.33167 71.08437 258.83167Q71.70937 258.33167 72.64687 258.33167Q73.33437 258.33167 73.85 258.6129Q74.38125 258.89417 74.725 259.4254Q75.06875 259.94104 75.25625 260.69104Q75.44375 261.44104 75.44375 262.7223Q75.44375 264.2379 75.13125 265.1598Q74.83437 266.08167 74.20937 266.5973Q73.6 267.0973 72.64687 267.0973Q71.4125 267.0973 70.69375 266.20667Q69.85 265.14417 69.85 262.7223ZM70.92812 262.7223Q70.92812 264.83167 71.42812 265.5348Q71.92812 266.2379 72.64687 266.2379Q73.38125 266.2379 73.86562 265.5348Q74.36562 264.83167 74.36562 262.7223Q74.36562 260.5973 73.86562 259.9098Q73.38125 259.20667 72.63125 259.20667Q71.9125 259.20667 71.475 259.81604Q70.92812 260.5973 70.92812 262.7223Z" fill-rule="nonzero"/><path fill="#000000" d="M55.85 200.34438L56.9125 200.20375Q57.084373 201.09438 57.521873 201.50063Q57.959373 201.89125 58.584373 201.89125Q59.334373 201.89125 59.85 201.37563Q60.365623 200.86 60.365623 200.09438Q60.365623 199.37563 59.88125 198.90688Q59.4125 198.4225 58.678123 198.4225Q58.38125 198.4225 57.928123 198.5475L58.053123 197.61Q58.1625 197.62563 58.225 197.62563Q58.896873 197.62563 59.428123 197.28188Q59.975 196.9225 59.975 196.18813Q59.975 195.61 59.584373 195.235Q59.19375 194.84438 58.56875 194.84438Q57.959373 194.84438 57.5375 195.235Q57.13125 195.62563 57.021873 196.39125L55.959373 196.20375Q56.146873 195.15688 56.834373 194.57875Q57.521873 193.985 58.5375 193.985Q59.240623 193.985 59.834373 194.2975Q60.428123 194.59438 60.740623 195.11Q61.053123 195.62563 61.053123 196.21938Q61.053123 196.76625 60.75625 197.235Q60.459373 197.68813 59.88125 197.95375Q60.63125 198.12563 61.053123 198.68813Q61.475 199.235 61.475 200.07875Q61.475 201.20375 60.646873 201.985Q59.834373 202.76625 58.584373 202.76625Q57.459373 202.76625 56.709373 202.09438Q55.959373 201.40688 55.85 200.34438ZM62.85 198.37563Q62.85 196.84438 63.1625 195.9225Q63.475 194.985 64.08437 194.485Q64.70937 193.985 65.64687 193.985Q66.33437 193.985 66.85 194.26625Q67.38125 194.5475 67.725 195.07875Q68.06875 195.59438 68.25625 196.34438Q68.44375 197.09438 68.44375 198.37563Q68.44375 199.89125 68.13125 200.81313Q67.83437 201.735 67.20937 202.25063Q66.6 202.75063 65.64687 202.75063Q64.4125 202.75063 63.69375 201.86Q62.85 200.7975 62.85 198.37563ZM63.928123 198.37563Q63.928123 200.485 64.42812 201.18813Q64.92812 201.89125 65.64687 201.89125Q66.38125 201.89125 66.86562 201.18813Q67.36562 200.485 67.36562 198.37563Q67.36562 196.25063 66.86562 195.56313Q66.38125 194.86 65.63125 194.86Q64.9125 194.86 64.475 195.46938Q63.928123 196.25063 63.928123 198.37563ZM69.85 198.37563Q69.85 196.84438 70.1625 195.9225Q70.475 194.985 71.08437 194.485Q71.70937 193.985 72.64687 193.985Q73.33437 193.985 73.85 194.26625Q74.38125 194.5475 74.725 195.07875Q75.06875 195.59438 75.25625 196.34438Q75.44375 197.09438 75.44375 198.37563Q75.44375 199.89125 75.13125 200.81313Q74.83437 201.735 74.20937 202.25063Q73.6 202.75063 72.64687 202.75063Q71.4125 202.75063 70.69375 201.86Q69.85 200.7975 69.85 198.37563ZM70.92812 198.37563Q70.92812 200.485 71.42812 201.18813Q71.92812 201.89125 72.64687 201.89125Q73.38125 201.89125 73.86562 201.18813Q74.36562 200.485 74.36562 198.37563Q74.36562 196.25063 73.86562 195.56313Q73.38125 194.86 72.63125 194.86Q71.9125 194.86 71.475 195.46938Q70.92812 196.25063 70.92812 198.37563Z" fill-rule="nonzero"/><path fill="#000000" d="M59.225 138.26334L59.225 136.20084L55.50625 136.20084L55.50625 135.23209L59.428123 129.66959L60.2875 129.66959L60.2875 135.23209L61.44375 135.23209L61.44375 136.20084L60.2875 136.20084L60.2875 138.26334L59.225 138.26334ZM59.225 135.23209L59.225 131.37271L56.5375 135.23209L59.225 135.23209ZM62.85 134.02896Q62.85 132.49771 63.1625 131.57584Q63.475 130.63834 64.08437 130.13834Q64.70937 129.63834 65.64687 129.63834Q66.33437 129.63834 66.85 129.91959Q67.38125 130.20084 67.725 130.73209Q68.06875 131.24771 68.25625 131.99771Q68.44375 132.74771 68.44375 134.02896Q68.44375 135.54459 68.13125 136.46646Q67.83437 137.38834 67.20937 137.90396Q66.6 138.40396 65.64687 138.40396Q64.4125 138.40396 63.69375 137.51334Q62.85 136.45084 62.85 134.02896ZM63.928123 134.02896Q63.928123 136.13834 64.42812 136.84146Q64.92812 137.54459 65.64687 137.54459Q66.38125 137.54459 66.86562 136.84146Q67.36562 136.13834 67.36562 134.02896Q67.36562 131.90396 66.86562 131.21646Q66.38125 130.51334 65.63125 130.51334Q64.9125 130.51334 64.475 131.12271Q63.928123 131.90396 63.928123 134.02896ZM69.85 134.02896Q69.85 132.49771 70.1625 131.57584Q70.475 130.63834 71.08437 130.13834Q71.70937 129.63834 72.64687 129.63834Q73.33437 129.63834 73.85 129.91959Q74.38125 130.20084 74.725 130.73209Q75.06875 131.24771 75.25625 131.99771Q75.44375 132.74771 75.44375 134.02896Q75.44375 135.54459 75.13125 136.46646Q74.83437 137.38834 74.20937 137.90396Q73.6 138.40396 72.64687 138.40396Q71.4125 138.40396 70.69375 137.51334Q69.85 136.45084 69.85 134.02896ZM70.92812 134.02896Q70.92812 136.13834 71.42812 136.84146Q71.92812 137.54459 72.64687 137.54459Q73.38125 137.54459 73.86562 136.84146Q74.36562 136.13834 74.36562 134.02896Q74.36562 131.90396 73.86562 131.21646Q73.38125 130.51334 72.63125 130.51334Q71.9125 130.51334 71.475 131.12271Q70.92812 131.90396 70.92812 134.02896Z" fill-rule="nonzero"/><path fill="#000000" d="M55.85 71.666664L56.959373 71.572914Q57.084373 72.385414 57.521873 72.791664Q57.975 73.197914 58.615623 73.197914Q59.365623 73.197914 59.896873 72.61979Q60.428123 72.041664 60.428123 71.104164Q60.428123 70.197914 59.9125 69.68229Q59.4125 69.15104 58.584373 69.15104Q58.084373 69.15104 57.6625 69.385414Q57.25625 69.61979 57.021873 69.979164L56.0375 69.854164L56.865623 65.447914L61.146873 65.447914L61.146873 66.447914L57.709373 66.447914L57.25625 68.760414Q58.021873 68.21354 58.865623 68.21354Q59.990623 68.21354 60.75625 68.99479Q61.5375 69.77604 61.5375 71.010414Q61.5375 72.166664 60.865623 73.02604Q60.0375 74.05729 58.615623 74.05729Q57.44375 74.05729 56.69375 73.40104Q55.959373 72.74479 55.85 71.666664ZM62.85 69.68229Q62.85 68.15104 63.1625 67.229164Q63.475 66.291664 64.08437 65.791664Q64.70937 65.291664 65.64687 65.291664Q66.33437 65.291664 66.85 65.572914Q67.38125 65.854164 67.725 66.385414Q68.06875 66.90104 68.25625 67.65104Q68.44375 68.40104 68.44375 69.68229Q68.44375 71.197914 68.13125 72.11979Q67.83437 73.041664 67.20937 73.55729Q66.6 74.05729 65.64687 74.05729Q64.4125 74.05729 63.69375 73.166664Q62.85 72.104164 62.85 69.68229ZM63.928123 69.68229Q63.928123 71.791664 64.42812 72.49479Q64.92812 73.197914 65.64687 73.197914Q66.38125 73.197914 66.86562 72.49479Q67.36562 71.791664 67.36562 69.68229Q67.36562 67.55729 66.86562 66.86979Q66.38125 66.166664 65.63125 66.166664Q64.9125 66.166664 64.475 66.77604Q63.928123 67.55729 63.928123 69.68229ZM69.85 69.68229Q69.85 68.15104 70.1625 67.229164Q70.475 66.291664 71.08437 65.791664Q71.70937 65.291664 72.64687 65.291664Q73.33437 65.291664 73.85 65.572914Q74.38125 65.854164 74.725 66.385414Q75.06875 66.90104 75.25625 67.65104Q75.44375 68.40104 75.44375 69.68229Q75.44375 71.197914 75.13125 72.11979Q74.83437 73.041664 74.20937 73.55729Q73.6 74.05729 72.64687 74.05729Q71.4125 74.05729 70.69375 73.166664Q69.85 72.104164 69.85 69.68229ZM70.92812 69.68229Q70.92812 71.791664 71.42812 72.49479Q71.92812 73.197914 72.64687 73.197914Q73.38125 73.197914 73.86562 72.49479Q74.36562 71.791664 74.36562 69.68229Q74.36562 67.55729 73.86562 66.86979Q73.38125 66.166664 72.63125 66.166664Q71.9125 66.166664 71.475 66.77604Q70.92812 67.55729 70.92812 69.68229Z" fill-rule="nonzero"/><path fill="#000000" d="M90.40883 406.07187L90.40883 405.00937L93.65883 405.00937L93.65883 406.07187L90.40883 406.07187Z" fill-rule="nonzero"/><path fill="#000000" d="M144.31754 406.4L145.42691 406.30624Q145.55191 407.11874 145.98941 407.525Q146.44254 407.93124 147.08316 407.93124Q147.83316 407.93124 148.36441 407.35312Q148.89566 406.775 148.89566 405.8375Q148.89566 404.93124 148.38004 404.41562Q147.88004 403.88437 147.05191 403.88437Q146.55191 403.88437 146.13004 404.11874Q145.72379 404.35312 145.48941 404.7125L144.50504 404.5875L145.33316 400.18124L149.61441 400.18124L149.61441 401.18124L146.17691 401.18124L145.72379 403.49374Q146.48941 402.94687 147.33316 402.94687Q148.45816 402.94687 149.22379 403.72812Q150.00504 404.50937 150.00504 405.74374Q150.00504 406.9 149.33316 407.75937Q148.50504 408.79062 147.08316 408.79062Q145.91129 408.79062 145.16129 408.13437Q144.42691 407.47812 144.31754 406.4ZM151.31754 404.41562Q151.31754 402.88437 151.63004 401.9625Q151.94254 401.025 152.55191 400.525Q153.17691 400.025 154.11441 400.025Q154.80191 400.025 155.31754 400.30624Q155.84879 400.5875 156.19254 401.11874Q156.53629 401.63437 156.72379 402.38437Q156.91129 403.13437 156.91129 404.41562Q156.91129 405.93124 156.59879 406.85312Q156.30191 407.775 155.67691 408.29062Q155.06754 408.79062 154.11441 408.79062Q152.88004 408.79062 152.16129 407.9Q151.31754 406.8375 151.31754 404.41562ZM152.39566 404.41562Q152.39566 406.525 152.89566 407.22812Q153.39566 407.93124 154.11441 407.93124Q154.84879 407.93124 155.33316 407.22812Q155.83316 406.525 155.83316 404.41562Q155.83316 402.29062 155.33316 401.60312Q154.84879 400.9 154.09879 400.9Q153.38004 400.9 152.94254 401.50937Q152.39566 402.29062 152.39566 404.41562ZM158.88004 408.65L158.88004 407.44687L160.08316 407.44687L160.08316 408.65Q160.08316 409.30624 159.84879 409.7125Q159.61441 410.13437 159.11441 410.35312L158.81754 409.9Q159.14566 409.75937 159.30191 409.4625Q159.45816 409.18124 159.48941 408.65L158.88004 408.65ZM161.31754 404.41562Q161.31754 402.88437 161.63004 401.9625Q161.94254 401.025 162.55191 400.525Q163.17691 400.025 164.11441 400.025Q164.80191 400.025 165.31754 400.30624Q165.84879 400.5875 166.19254 401.11874Q166.53629 401.63437 166.72379 402.38437Q166.91129 403.13437 166.91129 404.41562Q166.91129 405.93124 166.59879 406.85312Q166.30191 407.775 165.67691 408.29062Q165.06754 408.79062 164.11441 408.79062Q162.88004 408.79062 162.16129 407.9Q161.31754 406.8375 161.31754 404.41562ZM162.39566 404.41562Q162.39566 406.525 162.89566 407.22812Q163.39566 407.93124 164.11441 407.93124Q164.84879 407.93124 165.33316 407.22812Q165.83316 406.525 165.83316 404.41562Q165.83316 402.29062 165.33316 401.60312Q164.84879 400.9 164.09879 400.9Q163.38004 400.9 162.94254 401.50937Q162.39566 402.29062 162.39566 404.41562ZM168.31754 404.41562Q168.31754 402.88437 168.63004 401.9625Q168.94254 401.025 169.55191 400.525Q170.17691 400.025 171.11441 400.025Q171.80191 400.025 172.31754 400.30624Q172.84879 400.5875 173.19254 401.11874Q173.53629 401.63437 173.72379 402.38437Q173.91129 403.13437 173.91129 404.41562Q173.91129 405.93124 173.59879 406.85312Q173.30191 407.775 172.67691 408.29062Q172.06754 408.79062 171.11441 408.79062Q169.88004 408.79062 169.16129 407.9Q168.31754 406.8375 168.31754 404.41562ZM169.39566 404.41562Q169.39566 406.525 169.89566 407.22812Q170.39566 407.93124 171.11441 407.93124Q171.84879 407.93124 172.33316 407.22812Q172.83316 406.525 172.83316 404.41562Q172.83316 402.29062 172.33316 401.60312Q171.84879 400.9 171.09879 400.9Q170.38004 400.9 169.94254 401.50937Q169.39566 402.29062 169.39566 404.41562ZM175.31754 404.41562Q175.31754 402.88437 175.63004 401.9625Q175.94254 401.025 176.55191 400.525Q177.17691 400.025 178.11441 400.025Q178.80191 400.025 179.31754 400.30624Q179.84879 400.5875 180.19254 401.11874Q180.53629 401.63437 180.72379 402.38437Q180.91129 403.13437 180.91129 404.41562Q180.91129 405.93124 180.59879 406.85312Q180.30191 407.775 179.67691 408.29062Q179.06754 408.79062 178.11441 408.79062Q176.88004 408.79062 176.16129 407.9Q175.31754 406.8375 175.31754 404.41562ZM176.39566 404.41562Q176.39566 406.525 176.89566 407.22812Q177.39566 407.93124 178.11441 407.93124Q178.84879 407.93124 179.33316 407.22812Q179.83316 406.525 179.83316 404.41562Q179.83316 402.29062 179.33316 401.60312Q178.84879 400.9 178.09879 400.9Q177.38004 400.9 176.94254 401.50937Q176.39566 402.29062 176.39566 404.41562ZM182.88004 408.65L182.88004 407.44687L184.08316 407.44687L184.08316 408.65Q184.08316 409.30624 183.84879 409.7125Q183.61441 410.13437 183.11441 410.35312L182.81754 409.9Q183.14566 409.75937 183.30191 409.4625Q183.45816 409.18124 183.48941 408.65L182.88004 408.65ZM185.31754 404.41562Q185.31754 402.88437 185.63004 401.9625Q185.94254 401.025 186.55191 400.525Q187.17691 400.025 188.11441 400.025Q188.80191 400.025 189.31754 400.30624Q189.84879 400.5875 190.19254 401.11874Q190.53629 401.63437 190.72379 402.38437Q190.91129 403.13437 190.91129 404.41562Q190.91129 405.93124 190.59879 406.85312Q190.30191 407.775 189.67691 408.29062Q189.06754 408.79062 188.11441 408.79062Q186.88004 408.79062 186.16129 407.9Q185.31754 406.8375 185.31754 404.41562ZM186.39566 404.41562Q186.39566 406.525 186.89566 407.22812Q187.39566 407.93124 188.11441 407.93124Q188.84879 407.93124 189.33316 407.22812Q189.83316 406.525 189.83316 404.41562Q189.83316 402.29062 189.33316 401.60312Q188.84879 400.9 188.09879 400.9Q187.38004 400.9 186.94254 401.50937Q186.39566 402.29062 186.39566 404.41562ZM192.31754 404.41562Q192.31754 402.88437 192.63004 401.9625Q192.94254 401.025 193.55191 400.525Q194.17691 400.025 195.11441 400.025Q195.80191 400.025 196.31754 400.30624Q196.84879 400.5875 197.19254 401.11874Q197.53629 401.63437 197.72379 402.38437Q197.91129 403.13437 197.91129 404.41562Q197.91129 405.93124 197.59879 406.85312Q197.30191 407.775 196.67691 408.29062Q196.06754 408.79062 195.11441 408.79062Q193.88004 408.79062 193.16129 407.9Q192.31754 406.8375 192.31754 404.41562ZM193.39566 404.41562Q193.39566 406.525 193.89566 407.22812Q194.39566 407.93124 195.11441 407.93124Q195.84879 407.93124 196.33316 407.22812Q196.83316 406.525 196.83316 404.41562Q196.83316 402.29062 196.33316 401.60312Q195.84879 400.9 195.09879 400.9Q194.38004 400.9 193.94254 401.50937Q193.39566 402.29062 193.39566 404.41562ZM199.31754 404.41562Q199.31754 402.88437 199.63004 401.9625Q199.94254 401.025 200.55191 400.525Q201.17691 400.025 202.11441 400.025Q202.80191 400.025 203.31754 400.30624Q203.84879 400.5875 204.19254 401.11874Q204.53629 401.63437 204.72379 402.38437Q204.91129 403.13437 204.91129 404.41562Q204.91129 405.93124 204.59879 406.85312Q204.30191 407.775 203.67691 408.29062Q203.06754 408.79062 202.11441 408.79062Q200.88004 408.79062 200.16129 407.9Q199.31754 406.8375 199.31754 404.41562ZM200.39566 404.41562Q200.39566 406.525 200.89566 407.22812Q201.39566 407.93124 202.11441 407.93124Q202.84879 407.93124 203.33316 407.22812Q203.83316 406.525 203.83316 404.41562Q203.83316 402.29062 203.33316 401.60312Q202.84879 400.9 202.09879 400.9Q201.38004 400.9 200.94254 401.50937Q200.39566 402.29062 200.39566 404.41562Z" fill-rule="nonzero"/><path fill="#000000" d="M227.56998 408.65L226.5231 408.65L226.5231 401.93124Q226.13248 402.29062 225.50748 402.66562Q224.8981 403.025 224.41373 403.19687L224.41373 402.18124Q225.28873 401.75937 225.94498 401.18124Q226.61685 400.5875 226.8981 400.025L227.56998 400.025L227.56998 408.65ZM230.60123 404.41562Q230.60123 402.88437 230.91373 401.9625Q231.22623 401.025 231.8356 400.525Q232.4606 400.025 233.3981 400.025Q234.0856 400.025 234.60123 400.30624Q235.13248 400.5875 235.47623 401.11874Q235.81998 401.63437 236.00748 402.38437Q236.19498 403.13437 236.19498 404.41562Q236.19498 405.93124 235.88248 406.85312Q235.5856 407.775 234.9606 408.29062Q234.35123 408.79062 233.3981 408.79062Q232.16373 408.79062 231.44498 407.9Q230.60123 406.8375 230.60123 404.41562ZM231.67935 404.41562Q231.67935 406.525 232.17935 407.22812Q232.67935 407.93124 233.3981 407.93124Q234.13248 407.93124 234.61685 407.22812Q235.11685 406.525 235.11685 404.41562Q235.11685 402.29062 234.61685 401.60312Q234.13248 400.9 233.38248 400.9Q232.66373 400.9 232.22623 401.50937Q231.67935 402.29062 231.67935 404.41562ZM237.60123 404.41562Q237.60123 402.88437 237.91373 401.9625Q238.22623 401.025 238.8356 400.525Q239.4606 400.025 240.3981 400.025Q241.0856 400.025 241.60123 400.30624Q242.13248 400.5875 242.47623 401.11874Q242.81998 401.63437 243.00748 402.38437Q243.19498 403.13437 243.19498 404.41562Q243.19498 405.93124 242.88248 406.85312Q242.5856 407.775 241.9606 408.29062Q241.35123 408.79062 240.3981 408.79062Q239.16373 408.79062 238.44498 407.9Q237.60123 406.8375 237.60123 404.41562ZM238.67935 404.41562Q238.67935 406.525 239.17935 407.22812Q239.67935 407.93124 240.3981 407.93124Q241.13248 407.93124 241.61685 407.22812Q242.11685 406.525 242.11685 404.41562Q242.11685 402.29062 241.61685 401.60312Q241.13248 400.9 240.38248 400.9Q239.66373 400.9 239.22623 401.50937Q238.67935 402.29062 238.67935 404.41562ZM245.16373 408.65L245.16373 407.44687L246.36685 407.44687L246.36685 408.65Q246.36685 409.30624 246.13248 409.7125Q245.8981 410.13437 245.3981 410.35312L245.10123 409.9Q245.42935 409.75937 245.5856 409.4625Q245.74185 409.18124 245.7731 408.65L245.16373 408.65ZM247.60123 404.41562Q247.60123 402.88437 247.91373 401.9625Q248.22623 401.025 248.8356 400.525Q249.4606 400.025 250.3981 400.025Q251.0856 400.025 251.60123 400.30624Q252.13248 400.5875 252.47623 401.11874Q252.81998 401.63437 253.00748 402.38437Q253.19498 403.13437 253.19498 404.41562Q253.19498 405.93124 252.88248 406.85312Q252.5856 407.775 251.9606 408.29062Q251.35123 408.79062 250.3981 408.79062Q249.16373 408.79062 248.44498 407.9Q247.60123 406.8375 247.60123 404.41562ZM248.67935 404.41562Q248.67935 406.525 249.17935 407.22812Q249.67935 407.93124 250.3981 407.93124Q251.13248 407.93124 251.61685 407.22812Q252.11685 406.525 252.11685 404.41562Q252.11685 402.29062 251.61685 401.60312Q251.13248 400.9 250.38248 400.9Q249.66373 400.9 249.22623 401.50937Q248.67935 402.29062 248.67935 404.41562ZM254.60123 404.41562Q254.60123 402.88437 254.91373 401.9625Q255.22623 401.025 255.8356 400.525Q256.4606 400.025 257.3981 400.025Q258.0856 400.025 258.60123 400.30624Q259.13248 400.5875 259.47623 401.11874Q259.81998 401.63437 260.00748 402.38437Q260.19498 403.13437 260.19498 404.41562Q260.19498 405.93124 259.88248 406.85312Q259.5856 407.775 258.9606 408.29062Q258.35123 408.79062 257.3981 408.79062Q256.16373 408.79062 255.44498 407.9Q254.60123 406.8375 254.60123 404.41562ZM255.67935 404.41562Q255.67935 406.525 256.17935 407.22812Q256.67935 407.93124 257.3981 407.93124Q258.13248 407.93124 258.61685 407.22812Q259.11685 406.525 259.11685 404.41562Q259.11685 402.29062 258.61685 401.60312Q258.13248 400.9 257.38248 400.9Q256.66373 400.9 256.22623 401.50937Q255.67935 402.29062 255.67935 404.41562ZM261.60123 404.41562Q261.60123 402.88437 261.91373 401.9625Q262.22623 401.025 262.8356 400.525Q263.4606 400.025 264.3981 400.025Q265.0856 400.025 265.60123 400.30624Q266.13248 400.5875 266.47623 401.11874Q266.81998 401.63437 267.00748 402.38437Q267.19498 403.13437 267.19498 404.41562Q267.19498 405.93124 266.88248 406.85312Q266.5856 407.775 265.9606 408.29062Q265.35123 408.79062 264.3981 408.79062Q263.16373 408.79062 262.44498 407.9Q261.60123 406.8375 261.60123 404.41562ZM262.67935 404.41562Q262.67935 406.525 263.17935 407.22812Q263.67935 407.93124 264.3981 407.93124Q265.13248 407.93124 265.61685 407.22812Q266.11685 406.525 266.11685 404.41562Q266.11685 402.29062 265.61685 401.60312Q265.13248 400.9 264.38248 400.9Q263.66373 400.9 263.22623 401.50937Q262.67935 402.29062 262.67935 404.41562ZM269.16373 408.65L269.16373 407.44687L270.36685 407.44687L270.36685 408.65Q270.36685 409.30624 270.13248 409.7125Q269.8981 410.13437 269.3981 410.35312L269.10123 409.9Q269.42935 409.75937 269.5856 409.4625Q269.74185 409.18124 269.7731 408.65L269.16373 408.65ZM271.60123 404.41562Q271.60123 402.88437 271.91373 401.9625Q272.22623 401.025 272.8356 400.525Q273.4606 400.025 274.3981 400.025Q275.0856 400.025 275.60123 400.30624Q276.13248 400.5875 276.47623 401.11874Q276.81998 401.63437 277.00748 402.38437Q277.19498 403.13437 277.19498 404.41562Q277.19498 405.93124 276.88248 406.85312Q276.5856 407.775 275.9606 408.29062Q275.35123 408.79062 274.3981 408.79062Q273.16373 408.79062 272.44498 407.9Q271.60123 406.8375 271.60123 404.41562ZM272.67935 404.41562Q272.67935 406.525 273.17935 407.22812Q273.67935 407.93124 274.3981 407.93124Q275.13248 407.93124 275.61685 407.22812Q276.11685 406.525 276.11685 404.41562Q276.11685 402.29062 275.61685 401.60312Q275.13248 400.9 274.38248 400.9Q273.66373 400.9 273.22623 401.50937Q272.67935 402.29062 272.67935 404.41562ZM278.60123 404.41562Q278.60123 402.88437 278.91373 401.9625Q279.22623 401.025 279.8356 400.525Q280.4606 400.025 281.3981 400.025Q282.0856 400.025 282.60123 400.30624Q283.13248 400.5875 283.47623 401.11874Q283.81998 401.63437 284.00748 402.38437Q284.19498 403.13437 284.19498 404.41562Q284.19498 405.93124 283.88248 406.85312Q283.5856 407.775 282.9606 408.29062Q282.35123 408.79062 281.3981 408.79062Q280.16373 408.79062 279.44498 407.9Q278.60123 406.8375 278.60123 404.41562ZM279.67935 404.41562Q279.67935 406.525 280.17935 407.22812Q280.67935 407.93124 281.3981 407.93124Q282.13248 407.93124 282.61685 407.22812Q283.11685 406.525 283.11685 404.41562Q283.11685 402.29062 282.61685 401.60312Q282.13248 400.9 281.38248 400.9Q280.66373 400.9 280.22623 401.50937Q279.67935 402.29062 279.67935 404.41562ZM285.60123 404.41562Q285.60123 402.88437 285.91373 401.9625Q286.22623 401.025 286.8356 400.525Q287.4606 400.025 288.3981 400.025Q289.0856 400.025 289.60123 400.30624Q290.13248 400.5875 290.47623 401.11874Q290.81998 401.63437 291.00748 402.38437Q291.19498 403.13437 291.19498 404.41562Q291.19498 405.93124 290.88248 406.85312Q290.5856 407.775 289.9606 408.29062Q289.35123 408.79062 288.3981 408.79062Q287.16373 408.79062 286.44498 407.9Q285.60123 406.8375 285.60123 404.41562ZM286.67935 404.41562Q286.67935 406.525 287.17935 407.22812Q287.67935 407.93124 288.3981 407.93124Q289.13248 407.93124 289.61685 407.22812Q290.11685 406.525 290.11685 404.41562Q290.11685 402.29062 289.61685 401.60312Q289.13248 400.9 288.38248 400.9Q287.66373 400.9 287.22623 401.50937Q286.67935 402.29062 286.67935 404.41562Z" fill-rule="nonzero"/><path fill="#000000" d="M310.3537 408.65L309.30682 408.65L309.30682 401.93124Q308.9162 402.29062 308.2912 402.66562Q307.68182 403.025 307.19745 403.19687L307.19745 402.18124Q308.07245 401.75937 308.7287 401.18124Q309.40057 400.5875 309.68182 400.025L310.3537 400.025L310.3537 408.65ZM313.38495 406.4L314.49432 406.30624Q314.61932 407.11874 315.05682 407.525Q315.50995 407.93124 316.15057 407.93124Q316.90057 407.93124 317.43182 407.35312Q317.96307 406.775 317.96307 405.8375Q317.96307 404.93124 317.44745 404.41562Q316.94745 403.88437 316.11932 403.88437Q315.61932 403.88437 315.19745 404.11874Q314.7912 404.35312 314.55682 404.7125L313.57245 404.5875L314.40057 400.18124L318.68182 400.18124L318.68182 401.18124L315.24432 401.18124L314.7912 403.49374Q315.55682 402.94687 316.40057 402.94687Q317.52557 402.94687 318.2912 403.72812Q319.07245 404.50937 319.07245 405.74374Q319.07245 406.9 318.40057 407.75937Q317.57245 408.79062 316.15057 408.79062Q314.9787 408.79062 314.2287 408.13437Q313.49432 407.47812 313.38495 406.4ZM320.38495 404.41562Q320.38495 402.88437 320.69745 401.9625Q321.00995 401.025 321.61932 400.525Q322.24432 400.025 323.18182 400.025Q323.86932 400.025 324.38495 400.30624Q324.9162 400.5875 325.25995 401.11874Q325.6037 401.63437 325.7912 402.38437Q325.9787 403.13437 325.9787 404.41562Q325.9787 405.93124 325.6662 406.85312Q325.36932 407.775 324.74432 408.29062Q324.13495 408.79062 323.18182 408.79062Q321.94745 408.79062 321.2287 407.9Q320.38495 406.8375 320.38495 404.41562ZM321.46307 404.41562Q321.46307 406.525 321.96307 407.22812Q322.46307 407.93124 323.18182 407.93124Q323.9162 407.93124 324.40057 407.22812Q324.90057 406.525 324.90057 404.41562Q324.90057 402.29062 324.40057 401.60312Q323.9162 400.9 323.1662 400.9Q322.44745 400.9 322.00995 401.50937Q321.46307 402.29062 321.46307 404.41562ZM327.94745 408.65L327.94745 407.44687L329.15057 407.44687L329.15057 408.65Q329.15057 409.30624 328.9162 409.7125Q328.68182 410.13437 328.18182 410.35312L327.88495 409.9Q328.21307 409.75937 328.36932 409.4625Q328.52557 409.18124 328.55682 408.65L327.94745 408.65ZM330.38495 404.41562Q330.38495 402.88437 330.69745 401.9625Q331.00995 401.025 331.61932 400.525Q332.24432 400.025 333.18182 400.025Q333.86932 400.025 334.38495 400.30624Q334.9162 400.5875 335.25995 401.11874Q335.6037 401.63437 335.7912 402.38437Q335.9787 403.13437 335.9787 404.41562Q335.9787 405.93124 335.6662 406.85312Q335.36932 407.775 334.74432 408.29062Q334.13495 408.79062 333.18182 408.79062Q331.94745 408.79062 331.2287 407.9Q330.38495 406.8375 330.38495 404.41562ZM331.46307 404.41562Q331.46307 406.525 331.96307 407.22812Q332.46307 407.93124 333.18182 407.93124Q333.9162 407.93124 334.40057 407.22812Q334.90057 406.525 334.90057 404.41562Q334.90057 402.29062 334.40057 401.60312Q333.9162 400.9 333.1662 400.9Q332.44745 400.9 332.00995 401.50937Q331.46307 402.29062 331.46307 404.41562ZM337.38495 404.41562Q337.38495 402.88437 337.69745 401.9625Q338.00995 401.025 338.61932 400.525Q339.24432 400.025 340.18182 400.025Q340.86932 400.025 341.38495 400.30624Q341.9162 400.5875 342.25995 401.11874Q342.6037 401.63437 342.7912 402.38437Q342.9787 403.13437 342.9787 404.41562Q342.9787 405.93124 342.6662 406.85312Q342.36932 407.775 341.74432 408.29062Q341.13495 408.79062 340.18182 408.79062Q338.94745 408.79062 338.2287 407.9Q337.38495 406.8375 337.38495 404.41562ZM338.46307 404.41562Q338.46307 406.525 338.96307 407.22812Q339.46307 407.93124 340.18182 407.93124Q340.9162 407.93124 341.40057 407.22812Q341.90057 406.525 341.90057 404.41562Q341.90057 402.29062 341.40057 401.60312Q340.9162 400.9 340.1662 400.9Q339.44745 400.9 339.00995 401.50937Q338.46307 402.29062 338.46307 404.41562ZM344.38495 404.41562Q344.38495 402.88437 344.69745 401.9625Q345.00995 401.025 345.61932 400.525Q346.24432 400.025 347.18182 400.025Q347.86932 400.025 348.38495 400.30624Q348.9162 400.5875 349.25995 401.11874Q349.6037 401.63437 349.7912 402.38437Q349.9787 403.13437 349.9787 404.41562Q349.9787 405.93124 349.6662 406.85312Q349.36932 407.775 348.74432 408.29062Q348.13495 408.79062 347.18182 408.79062Q345.94745 408.79062 345.2287 407.9Q344.38495 406.8375 344.38495 404.41562ZM345.46307 404.41562Q345.46307 406.525 345.96307 407.22812Q346.46307 407.93124 347.18182 407.93124Q347.9162 407.93124 348.40057 407.22812Q348.90057 406.525 348.90057 404.41562Q348.90057 402.29062 348.40057 401.60312Q347.9162 400.9 347.1662 400.9Q346.44745 400.9 346.00995 401.50937Q345.46307 402.29062 345.46307 404.41562ZM351.94745 408.65L351.94745 407.44687L353.15057 407.44687L353.15057 408.65Q353.15057 409.30624 352.9162 409.7125Q352.68182 410.13437 352.18182 410.35312L351.88495 409.9Q352.21307 409.75937 352.36932 409.4625Q352.52557 409.18124 352.55682 408.65L351.94745 408.65ZM354.38495 404.41562Q354.38495 402.88437 354.69745 401.9625Q355.00995 401.025 355.61932 400.525Q356.24432 400.025 357.18182 400.025Q357.86932 400.025 358.38495 400.30624Q358.9162 400.5875 359.25995 401.11874Q359.6037 401.63437 359.7912 402.38437Q359.9787 403.13437 359.9787 404.41562Q359.9787 405.93124 359.6662 406.85312Q359.36932 407.775 358.74432 408.29062Q358.13495 408.79062 357.18182 408.79062Q355.94745 408.79062 355.2287 407.9Q354.38495 406.8375 354.38495 404.41562ZM355.46307 404.41562Q355.46307 406.525 355.96307 407.22812Q356.46307 407.93124 357.18182 407.93124Q357.9162 407.93124 358.40057 407.22812Q358.90057 406.525 358.90057 404.41562Q358.90057 402.29062 358.40057 401.60312Q357.9162 400.9 357.1662 400.9Q356.44745 400.9 356.00995 401.50937Q355.46307 402.29062 355.46307 404.41562ZM361.38495 404.41562Q361.38495 402.88437 361.69745 401.9625Q362.00995 401.025 362.61932 400.525Q363.24432 400.025 364.18182 400.025Q364.86932 400.025 365.38495 400.30624Q365.9162 400.5875 366.25995 401.11874Q366.6037 401.63437 366.7912 402.38437Q366.9787 403.13437 366.9787 404.41562Q366.9787 405.93124 366.6662 406.85312Q366.36932 407.775 365.74432 408.29062Q365.13495 408.79062 364.18182 408.79062Q362.94745 408.79062 362.2287 407.9Q361.38495 406.8375 361.38495 404.41562ZM362.46307 404.41562Q362.46307 406.525 362.96307 407.22812Q363.46307 407.93124 364.18182 407.93124Q364.9162 407.93124 365.40057 407.22812Q365.90057 406.525 365.90057 404.41562Q365.90057 402.29062 365.40057 401.60312Q364.9162 400.9 364.1662 400.9Q363.44745 400.9 363.00995 401.50937Q362.46307 402.29062 362.46307 404.41562ZM368.38495 404.41562Q368.38495 402.88437 368.69745 401.9625Q369.00995 401.025 369.61932 400.525Q370.24432 400.025 371.18182 400.025Q371.86932 400.025 372.38495 400.30624Q372.9162 400.5875 373.25995 401.11874Q373.6037 401.63437 373.7912 402.38437Q373.9787 403.13437 373.9787 404.41562Q373.9787 405.93124 373.6662 406.85312Q373.36932 407.775 372.74432 408.29062Q372.13495 408.79062 371.18182 408.79062Q369.94745 408.79062 369.2287 407.9Q368.38495 406.8375 368.38495 404.41562ZM369.46307 404.41562Q369.46307 406.525 369.96307 407.22812Q370.46307 407.93124 371.18182 407.93124Q371.9162 407.93124 372.40057 407.22812Q372.90057 406.525 372.90057 404.41562Q372.90057 402.29062 372.40057 401.60312Q371.9162 400.9 371.1662 400.9Q370.44745 400.9 370.00995 401.50937Q369.46307 402.29062 369.46307 404.41562Z" fill-rule="nonzero"/><path fill="#000000" d="M394.7155 407.63437L394.7155 408.65L389.028 408.65Q389.028 408.275 389.153 407.91562Q389.37177 407.3375 389.8405 406.775Q390.3249 406.2125 391.23114 405.47812Q392.6374 404.32187 393.12177 403.65Q393.62177 402.97812 393.62177 402.38437Q393.62177 401.75937 393.16864 401.3375Q392.7155 400.9 391.99677 400.9Q391.23114 400.9 390.778 401.35312Q390.3249 401.80624 390.30927 402.61874L389.23114 402.50937Q389.3405 401.29062 390.05927 400.66562Q390.79364 400.025 392.028 400.025Q393.2624 400.025 393.98114 400.7125Q394.6999 401.4 394.6999 402.41562Q394.6999 402.93124 394.48114 403.43124Q394.278 403.91562 393.778 404.47812Q393.29364 405.025 392.16864 405.97812Q391.2155 406.775 390.93427 407.07187Q390.66864 407.35312 390.49677 407.63437L394.7155 407.63437ZM396.16864 404.41562Q396.16864 402.88437 396.48114 401.9625Q396.79364 401.025 397.403 400.525Q398.028 400.025 398.9655 400.025Q399.653 400.025 400.16864 400.30624Q400.6999 400.5875 401.04364 401.11874Q401.3874 401.63437 401.5749 402.38437Q401.7624 403.13437 401.7624 404.41562Q401.7624 405.93124 401.4499 406.85312Q401.153 407.775 400.528 408.29062Q399.91864 408.79062 398.9655 408.79062Q397.73114 408.79062 397.0124 407.9Q396.16864 406.8375 396.16864 404.41562ZM397.24677 404.41562Q397.24677 406.525 397.74677 407.22812Q398.24677 407.93124 398.9655 407.93124Q399.6999 407.93124 400.18427 407.22812Q400.68427 406.525 400.68427 404.41562Q400.68427 402.29062 400.18427 401.60312Q399.6999 400.9 398.9499 400.9Q398.23114 400.9 397.79364 401.50937Q397.24677 402.29062 397.24677 404.41562ZM403.16864 404.41562Q403.16864 402.88437 403.48114 401.9625Q403.79364 401.025 404.403 400.525Q405.028 400.025 405.9655 400.025Q406.653 400.025 407.16864 400.30624Q407.6999 400.5875 408.04364 401.11874Q408.3874 401.63437 408.5749 402.38437Q408.7624 403.13437 408.7624 404.41562Q408.7624 405.93124 408.4499 406.85312Q408.153 407.775 407.528 408.29062Q406.91864 408.79062 405.9655 408.79062Q404.73114 408.79062 404.0124 407.9Q403.16864 406.8375 403.16864 404.41562ZM404.24677 404.41562Q404.24677 406.525 404.74677 407.22812Q405.24677 407.93124 405.9655 407.93124Q406.6999 407.93124 407.18427 407.22812Q407.68427 406.525 407.68427 404.41562Q407.68427 402.29062 407.18427 401.60312Q406.6999 400.9 405.9499 400.9Q405.23114 400.9 404.79364 401.50937Q404.24677 402.29062 404.24677 404.41562ZM410.73114 408.65L410.73114 407.44687L411.93427 407.44687L411.93427 408.65Q411.93427 409.30624 411.6999 409.7125Q411.4655 410.13437 410.9655 410.35312L410.66864 409.9Q410.99677 409.75937 411.153 409.4625Q411.30927 409.18124 411.3405 408.65L410.73114 408.65ZM413.16864 404.41562Q413.16864 402.88437 413.48114 401.9625Q413.79364 401.025 414.403 400.525Q415.028 400.025 415.9655 400.025Q416.653 400.025 417.16864 400.30624Q417.6999 400.5875 418.04364 401.11874Q418.3874 401.63437 418.5749 402.38437Q418.7624 403.13437 418.7624 404.41562Q418.7624 405.93124 418.4499 406.85312Q418.153 407.775 417.528 408.29062Q416.91864 408.79062 415.9655 408.79062Q414.73114 408.79062 414.0124 407.9Q413.16864 406.8375 413.16864 404.41562ZM414.24677 404.41562Q414.24677 406.525 414.74677 407.22812Q415.24677 407.93124 415.9655 407.93124Q416.6999 407.93124 417.18427 407.22812Q417.68427 406.525 417.68427 404.41562Q417.68427 402.29062 417.18427 401.60312Q416.6999 400.9 415.9499 400.9Q415.23114 400.9 414.79364 401.50937Q414.24677 402.29062 414.24677 404.41562ZM420.16864 404.41562Q420.16864 402.88437 420.48114 401.9625Q420.79364 401.025 421.403 400.525Q422.028 400.025 422.9655 400.025Q423.653 400.025 424.16864 400.30624Q424.6999 400.5875 425.04364 401.11874Q425.3874 401.63437 425.5749 402.38437Q425.7624 403.13437 425.7624 404.41562Q425.7624 405.93124 425.4499 406.85312Q425.153 407.775 424.528 408.29062Q423.91864 408.79062 422.9655 408.79062Q421.73114 408.79062 421.0124 407.9Q420.16864 406.8375 420.16864 404.41562ZM421.24677 404.41562Q421.24677 406.525 421.74677 407.22812Q422.24677 407.93124 422.9655 407.93124Q423.6999 407.93124 424.18427 407.22812Q424.68427 406.525 424.68427 404.41562Q424.68427 402.29062 424.18427 401.60312Q423.6999 400.9 422.9499 400.9Q422.23114 400.9 421.79364 401.50937Q421.24677 402.29062 421.24677 404.41562ZM427.16864 404.41562Q427.16864 402.88437 427.48114 401.9625Q427.79364 401.025 428.403 400.525Q429.028 400.025 429.9655 400.025Q430.653 400.025 431.16864 400.30624Q431.6999 400.5875 432.04364 401.11874Q432.3874 401.63437 432.5749 402.38437Q432.7624 403.13437 432.7624 404.41562Q432.7624 405.93124 432.4499 406.85312Q432.153 407.775 431.528 408.29062Q430.91864 408.79062 429.9655 408.79062Q428.73114 408.79062 428.0124 407.9Q427.16864 406.8375 427.16864 404.41562ZM428.24677 404.41562Q428.24677 406.525 428.74677 407.22812Q429.24677 407.93124 429.9655 407.93124Q430.6999 407.93124 431.18427 407.22812Q431.68427 406.525 431.68427 404.41562Q431.68427 402.29062 431.18427 401.60312Q430.6999 400.9 429.9499 400.9Q429.23114 400.9 428.79364 401.50937Q428.24677 402.29062 428.24677 404.41562ZM434.73114 408.65L434.73114 407.44687L435.93427 407.44687L435.93427 408.65Q435.93427 409.30624 435.6999 409.7125Q435.4655 410.13437 434.9655 410.35312L434.66864 409.9Q434.99677 409.75937 435.153 409.4625Q435.30927 409.18124 435.3405 408.65L434.73114 408.65ZM437.16864 404.41562Q437.16864 402.88437 437.48114 401.9625Q437.79364 401.025 438.403 400.525Q439.028 400.025 439.9655 400.025Q440.653 400.025 441.16864 400.30624Q441.6999 400.5875 442.04364 401.11874Q442.3874 401.63437 442.5749 402.38437Q442.7624 403.13437 442.7624 404.41562Q442.7624 405.93124 442.4499 406.85312Q442.153 407.775 441.528 408.29062Q440.91864 408.79062 439.9655 408.79062Q438.73114 408.79062 438.0124 407.9Q437.16864 406.8375 437.16864 404.41562ZM438.24677 404.41562Q438.24677 406.525 438.74677 407.22812Q439.24677 407.93124 439.9655 407.93124Q440.6999 407.93124 441.18427 407.22812Q441.68427 406.525 441.68427 404.41562Q441.68427 402.29062 441.18427 401.60312Q440.6999 400.9 439.9499 400.9Q439.23114 400.9 438.79364 401.50937Q438.24677 402.29062 438.24677 404.41562ZM444.16864 404.41562Q444.16864 402.88437 444.48114 401.9625Q444.79364 401.025 445.403 400.525Q446.028 400.025 446.9655 400.025Q447.653 400.025 448.16864 400.30624Q448.6999 400.5875 449.04364 401.11874Q449.3874 401.63437 449.5749 402.38437Q449.7624 403.13437 449.7624 404.41562Q449.7624 405.93124 449.4499 406.85312Q449.153 407.775 448.528 408.29062Q447.91864 408.79062 446.9655 408.79062Q445.73114 408.79062 445.0124 407.9Q444.16864 406.8375 444.16864 404.41562ZM445.24677 404.41562Q445.24677 406.525 445.74677 407.22812Q446.24677 407.93124 446.9655 407.93124Q447.6999 407.93124 448.18427 407.22812Q448.68427 406.525 448.68427 404.41562Q448.68427 402.29062 448.18427 401.60312Q447.6999 400.9 446.9499 400.9Q446.23114 400.9 445.79364 401.50937Q445.24677 402.29062 445.24677 404.41562ZM451.16864 404.41562Q451.16864 402.88437 451.48114 401.9625Q451.79364 401.025 452.403 400.525Q453.028 400.025 453.9655 400.025Q454.653 400.025 455.16864 400.30624Q455.6999 400.5875 456.04364 401.11874Q456.3874 401.63437 456.5749 402.38437Q456.7624 403.13437 456.7624 404.41562Q456.7624 405.93124 456.4499 406.85312Q456.153 407.775 455.528 408.29062Q454.91864 408.79062 453.9655 408.79062Q452.73114 408.79062 452.0124 407.9Q451.16864 406.8375 451.16864 404.41562ZM452.24677 404.41562Q452.24677 406.525 452.74677 407.22812Q453.24677 407.93124 453.9655 407.93124Q454.6999 407.93124 455.18427 407.22812Q455.68427 406.525 455.68427 404.41562Q455.68427 402.29062 455.18427 401.60312Q454.6999 400.9 453.9499 400.9Q453.23114 400.9 452.79364 401.50937Q452.24677 402.29062 452.24677 404.41562Z" fill-rule="nonzero"/><path fill="#000000" d="M477.4992 407.63437L477.4992 408.65L471.8117 408.65Q471.8117 408.275 471.9367 407.91562Q472.15546 407.3375 472.6242 406.775Q473.10858 406.2125 474.01483 405.47812Q475.42108 404.32187 475.90546 403.65Q476.40546 402.97812 476.40546 402.38437Q476.40546 401.75937 475.95233 401.3375Q475.4992 400.9 474.78046 400.9Q474.01483 400.9 473.5617 401.35312Q473.10858 401.80624 473.09296 402.61874L472.01483 402.50937Q472.1242 401.29062 472.84296 400.66562Q473.57733 400.025 474.8117 400.025Q476.04608 400.025 476.76483 400.7125Q477.48358 401.4 477.48358 402.41562Q477.48358 402.93124 477.26483 403.43124Q477.0617 403.91562 476.5617 404.47812Q476.07733 405.025 474.95233 405.97812Q473.9992 406.775 473.71796 407.07187Q473.45233 407.35312 473.28046 407.63437L477.4992 407.63437ZM478.95233 406.4L480.0617 406.30624Q480.1867 407.11874 480.6242 407.525Q481.07733 407.93124 481.71796 407.93124Q482.46796 407.93124 482.9992 407.35312Q483.53046 406.775 483.53046 405.8375Q483.53046 404.93124 483.01483 404.41562Q482.51483 403.88437 481.6867 403.88437Q481.1867 403.88437 480.76483 404.11874Q480.35858 404.35312 480.1242 404.7125L479.13983 404.5875L479.96796 400.18124L484.2492 400.18124L484.2492 401.18124L480.8117 401.18124L480.35858 403.49374Q481.1242 402.94687 481.96796 402.94687Q483.09296 402.94687 483.85858 403.72812Q484.63983 404.50937 484.63983 405.74374Q484.63983 406.9 483.96796 407.75937Q483.13983 408.79062 481.71796 408.79062Q480.54608 408.79062 479.79608 408.13437Q479.0617 407.47812 478.95233 406.4ZM485.95233 404.41562Q485.95233 402.88437 486.26483 401.9625Q486.57733 401.025 487.1867 400.525Q487.8117 400.025 488.7492 400.025Q489.4367 400.025 489.95233 400.30624Q490.48358 400.5875 490.82733 401.11874Q491.17108 401.63437 491.35858 402.38437Q491.54608 403.13437 491.54608 404.41562Q491.54608 405.93124 491.23358 406.85312Q490.9367 407.775 490.3117 408.29062Q489.70233 408.79062 488.7492 408.79062Q487.51483 408.79062 486.79608 407.9Q485.95233 406.8375 485.95233 404.41562ZM487.03046 404.41562Q487.03046 406.525 487.53046 407.22812Q488.03046 407.93124 488.7492 407.93124Q489.48358 407.93124 489.96796 407.22812Q490.46796 406.525 490.46796 404.41562Q490.46796 402.29062 489.96796 401.60312Q489.48358 400.9 488.73358 400.9Q488.01483 400.9 487.57733 401.50937Q487.03046 402.29062 487.03046 404.41562ZM493.51483 408.65L493.51483 407.44687L494.71796 407.44687L494.71796 408.65Q494.71796 409.30624 494.48358 409.7125Q494.2492 410.13437 493.7492 410.35312L493.45233 409.9Q493.78046 409.75937 493.9367 409.4625Q494.09296 409.18124 494.1242 408.65L493.51483 408.65ZM495.95233 404.41562Q495.95233 402.88437 496.26483 401.9625Q496.57733 401.025 497.1867 400.525Q497.8117 400.025 498.7492 400.025Q499.4367 400.025 499.95233 400.30624Q500.48358 400.5875 500.82733 401.11874Q501.17108 401.63437 501.35858 402.38437Q501.54608 403.13437 501.54608 404.41562Q501.54608 405.93124 501.23358 406.85312Q500.9367 407.775 500.3117 408.29062Q499.70233 408.79062 498.7492 408.79062Q497.51483 408.79062 496.79608 407.9Q495.95233 406.8375 495.95233 404.41562ZM497.03046 404.41562Q497.03046 406.525 497.53046 407.22812Q498.03046 407.93124 498.7492 407.93124Q499.48358 407.93124 499.96796 407.22812Q500.46796 406.525 500.46796 404.41562Q500.46796 402.29062 499.96796 401.60312Q499.48358 400.9 498.73358 400.9Q498.01483 400.9 497.57733 401.50937Q497.03046 402.29062 497.03046 404.41562ZM502.95233 404.41562Q502.95233 402.88437 503.26483 401.9625Q503.57733 401.025 504.1867 400.525Q504.8117 400.025 505.7492 400.025Q506.4367 400.025 506.95233 400.30624Q507.48358 400.5875 507.82733 401.11874Q508.17108 401.63437 508.35858 402.38437Q508.54608 403.13437 508.54608 404.41562Q508.54608 405.93124 508.23358 406.85312Q507.9367 407.775 507.3117 408.29062Q506.70233 408.79062 505.7492 408.79062Q504.51483 408.79062 503.79608 407.9Q502.95233 406.8375 502.95233 404.41562ZM504.03046 404.41562Q504.03046 406.525 504.53046 407.22812Q505.03046 407.93124 505.7492 407.93124Q506.48358 407.93124 506.96796 407.22812Q507.46796 406.525 507.46796 404.41562Q507.46796 402.29062 506.96796 401.60312Q506.48358 400.9 505.73358 400.9Q505.01483 400.9 504.57733 401.50937Q504.03046 402.29062 504.03046 404.41562ZM509.95233 404.41562Q509.95233 402.88437 510.26483 401.9625Q510.57733 401.025 511.1867 400.525Q511.8117 400.025 512.7492 400.025Q513.4367 400.025 513.95233 400.30624Q514.4836 400.5875 514.82733 401.11874Q515.1711 401.63437 515.3586 402.38437Q515.5461 403.13437 515.5461 404.41562Q515.5461 405.93124 515.2336 406.85312Q514.9367 407.775 514.3117 408.29062Q513.70233 408.79062 512.7492 408.79062Q511.51483 408.79062 510.79608 407.9Q509.95233 406.8375 509.95233 404.41562ZM511.03046 404.41562Q511.03046 406.525 511.53046 407.22812Q512.03046 407.93124 512.7492 407.93124Q513.4836 407.93124 513.96796 407.22812Q514.46796 406.525 514.46796 404.41562Q514.46796 402.29062 513.96796 401.60312Q513.4836 400.9 512.7336 400.9Q512.01483 400.9 511.57733 401.50937Q511.03046 402.29062 511.03046 404.41562ZM517.51483 408.65L517.51483 407.44687L518.71796 407.44687L518.71796 408.65Q518.71796 409.30624 518.4836 409.7125Q518.2492 410.13437 517.7492 410.35312L517.45233 409.9Q517.78046 409.75937 517.9367 409.4625Q518.09296 409.18124 518.1242 408.65L517.51483 408.65ZM519.95233 404.41562Q519.95233 402.88437 520.26483 401.9625Q520.57733 401.025 521.1867 400.525Q521.8117 400.025 522.7492 400.025Q523.4367 400.025 523.95233 400.30624Q524.4836 400.5875 524.82733 401.11874Q525.1711 401.63437 525.3586 402.38437Q525.5461 403.13437 525.5461 404.41562Q525.5461 405.93124 525.2336 406.85312Q524.9367 407.775 524.3117 408.29062Q523.70233 408.79062 522.7492 408.79062Q521.51483 408.79062 520.7961 407.9Q519.95233 406.8375 519.95233 404.41562ZM521.03046 404.41562Q521.03046 406.525 521.53046 407.22812Q522.03046 407.93124 522.7492 407.93124Q523.4836 407.93124 523.96796 407.22812Q524.46796 406.525 524.46796 404.41562Q524.46796 402.29062 523.96796 401.60312Q523.4836 400.9 522.7336 400.9Q522.01483 400.9 521.57733 401.50937Q521.03046 402.29062 521.03046 404.41562ZM526.95233 404.41562Q526.95233 402.88437 527.26483 401.9625Q527.57733 401.025 528.1867 400.525Q528.8117 400.025 529.7492 400.025Q530.4367 400.025 530.95233 400.30624Q531.4836 400.5875 531.82733 401.11874Q532.1711 401.63437 532.3586 402.38437Q532.5461 403.13437 532.5461 404.41562Q532.5461 405.93124 532.2336 406.85312Q531.9367 407.775 531.3117 408.29062Q530.70233 408.79062 529.7492 408.79062Q528.51483 408.79062 527.7961 407.9Q526.95233 406.8375 526.95233 404.41562ZM528.03046 404.41562Q528.03046 406.525 528.53046 407.22812Q529.03046 407.93124 529.7492 407.93124Q530.4836 407.93124 530.96796 407.22812Q531.46796 406.525 531.46796 404.41562Q531.46796 402.29062 530.96796 401.60312Q530.4836 400.9 529.7336 400.9Q529.01483 400.9 528.57733 401.50937Q528.03046 402.29062 528.03046 404.41562ZM533.95233 404.41562Q533.95233 402.88437 534.26483 401.9625Q534.57733 401.025 535.1867 400.525Q535.8117 400.025 536.7492 400.025Q537.4367 400.025 537.95233 400.30624Q538.4836 400.5875 538.82733 401.11874Q539.1711 401.63437 539.3586 402.38437Q539.5461 403.13437 539.5461 404.41562Q539.5461 405.93124 539.2336 406.85312Q538.9367 407.775 538.3117 408.29062Q537.70233 408.79062 536.7492 408.79062Q535.51483 408.79062 534.7961 407.9Q533.95233 406.8375 533.95233 404.41562ZM535.03046 404.41562Q535.03046 406.525 535.53046 407.22812Q536.03046 407.93124 536.7492 407.93124Q537.4836 407.93124 537.96796 407.22812Q538.46796 406.525 538.46796 404.41562Q538.46796 402.29062 537.96796 401.60312Q537.4836 400.9 536.7336 400.9Q536.01483 400.9 535.57733 401.50937Q535.03046 402.29062 535.03046 404.41562Z" fill-rule="nonzero"/><path fill="#4285f4" d="M577.65 69.916664C577.65 73.23038 574.9637 75.916664 571.65 75.916664C568.3363 75.916664 565.65 73.23038 565.65 69.916664C565.65 66.60296 568.3363 63.916668 571.65 63.916668C574.9637 63.916668 577.65 66.60296 577.65 69.916664Z" fill-rule="nonzero"/><path fill="#1a1a1a" d="M592.7125 70.90104L593.8375 71.197914Q593.47815 72.58854 592.5563 73.322914Q591.6344 74.05729 590.29065 74.05729Q588.9 74.05729 588.025 73.49479Q587.15 72.93229 586.6969 71.86979Q586.2438 70.791664 586.2438 69.55729Q586.2438 68.21354 586.7594 67.21354Q587.275 66.21354 588.2125 65.697914Q589.16565 65.18229 590.3063 65.18229Q591.5875 65.18229 592.4625 65.83854Q593.35315 66.49479 593.6969 67.68229L592.5719 67.947914Q592.275 67.010414 591.6969 66.58854Q591.1344 66.15104 590.275 66.15104Q589.29065 66.15104 588.6188 66.635414Q587.9625 67.104164 587.6813 67.90104Q587.41565 68.697914 587.41565 69.55729Q587.41565 70.65104 587.72815 71.46354Q588.0563 72.27604 588.72815 72.68229Q589.4 73.08854 590.1969 73.08854Q591.15 73.08854 591.8063 72.541664Q592.47815 71.99479 592.7125 70.90104ZM595.5719 73.916664L595.5719 65.322914L598.8219 65.322914Q599.66565 65.322914 600.1188 65.40104Q600.7594 65.510414 601.1813 65.80729Q601.6188 66.104164 601.8688 66.635414Q602.1344 67.166664 602.1344 67.80729Q602.1344 68.90104 601.4313 69.666664Q600.7438 70.416664 598.91565 70.416664L596.7125 70.416664L596.7125 73.916664L595.5719 73.916664ZM596.7125 69.416664L598.9313 69.416664Q600.04065 69.416664 600.4938 69.010414Q600.9625 68.58854 600.9625 67.83854Q600.9625 67.30729 600.6813 66.93229Q600.41565 66.541664 599.97815 66.416664Q599.6813 66.33854 598.91565 66.33854L596.7125 66.33854L596.7125 69.416664ZM609.2125 65.322914L610.35315 65.322914L610.35315 70.291664Q610.35315 71.58854 610.0563 72.354164Q609.7594 73.104164 608.9938 73.58854Q608.22815 74.05729 606.9938 74.05729Q605.79065 74.05729 605.0094 73.65104Q604.2438 73.229164 603.91565 72.447914Q603.5875 71.65104 603.5875 70.291664L603.5875 65.322914L604.72815 65.322914L604.72815 70.27604Q604.72815 71.40104 604.9313 71.93229Q605.15 72.46354 605.65 72.760414Q606.16565 73.041664 606.9 73.041664Q608.15 73.041664 608.6813 72.479164Q609.2125 71.90104 609.2125 70.27604L609.2125 65.322914Z" fill-rule="nonzero"/><path fill="#ea4335" d="M577.65 93.916664C577.65 97.23038 574.9637 99.916664 571.65 99.916664C568.3363 99.916664 565.65 97.23038 565.65 93.916664C565.65 90.60296 568.3363 87.916664 571.65 87.916664C574.9637 87.916664 577.65 90.60296 577.65 93.916664Z" fill-rule="nonzero"/><path fill="#1a1a1a" d="M586.5719 97.916664L586.5719 89.322914L587.72815 89.322914L592.2438 96.072914L592.2438 89.322914L593.3375 89.322914L593.3375 97.916664L592.16565 97.916664L587.65 91.166664L587.65 97.916664L586.5719 97.916664ZM599.5094 97.15104Q598.91565 97.65104 598.3688 97.854164Q597.8375 98.05729 597.2125 98.05729Q596.1813 98.05729 595.6344 97.55729Q595.0875 97.05729 595.0875 96.27604Q595.0875 95.822914 595.29065 95.447914Q595.4938 95.05729 595.8375 94.83854Q596.1813 94.604164 596.60315 94.49479Q596.9 94.40104 597.54065 94.322914Q598.8063 94.18229 599.41565 93.96354Q599.41565 93.74479 599.41565 93.697914Q599.41565 93.041664 599.1188 92.77604Q598.7125 92.43229 597.91565 92.43229Q597.1813 92.43229 596.8219 92.697914Q596.4625 92.947914 596.29065 93.604164L595.2594 93.46354Q595.4 92.80729 595.72815 92.40104Q596.0563 91.99479 596.66565 91.77604Q597.275 91.55729 598.0719 91.55729Q598.8688 91.55729 599.3688 91.74479Q599.8688 91.93229 600.10315 92.21354Q600.3375 92.49479 600.4313 92.93229Q600.47815 93.197914 600.47815 93.90104L600.47815 95.30729Q600.47815 96.77604 600.54065 97.166664Q600.6188 97.55729 600.8219 97.916664L599.7125 97.916664Q599.5563 97.58854 599.5094 97.15104ZM599.41565 94.791664Q598.8375 95.02604 597.6969 95.197914Q597.04065 95.291664 596.775 95.416664Q596.5094 95.52604 596.35315 95.74479Q596.2125 95.96354 596.2125 96.24479Q596.2125 96.666664 596.525 96.947914Q596.85315 97.229164 597.4625 97.229164Q598.0719 97.229164 598.54065 96.96354Q599.025 96.697914 599.2438 96.229164Q599.41565 95.86979 599.41565 95.18229L599.41565 94.791664ZM602.4469 90.541664L602.4469 89.322914L603.5094 89.322914L603.5094 90.541664L602.4469 90.541664ZM602.4469 97.916664L602.4469 91.697914L603.5094 91.697914L603.5094 97.916664L602.4469 97.916664ZM607.16565 97.916664L604.8063 91.697914L605.91565 91.697914L607.2594 95.416664Q607.4625 96.02604 607.65 96.666664Q607.79065 96.18229 608.04065 95.49479L609.4313 91.697914L610.5094 91.697914L608.15 97.916664L607.16565 97.916664ZM614.6969 95.916664L615.79065 96.041664Q615.54065 96.99479 614.8375 97.52604Q614.1344 98.05729 613.0563 98.05729Q611.6969 98.05729 610.8844 97.21354Q610.0875 96.36979 610.0875 94.854164Q610.0875 93.291664 610.9 92.43229Q611.7125 91.55729 612.9938 91.55729Q614.2438 91.55729 615.025 92.40104Q615.8219 93.24479 615.8219 94.791664Q615.8219 94.885414 615.8219 95.072914L611.1813 95.072914Q611.2438 96.104164 611.7594 96.65104Q612.275 97.18229 613.0563 97.18229Q613.6344 97.18229 614.04065 96.885414Q614.4625 96.572914 614.6969 95.916664ZM611.2438 94.21354L614.7125 94.21354Q614.65 93.416664 614.3219 93.02604Q613.8063 92.416664 613.0094 92.416664Q612.275 92.416664 611.775 92.90104Q611.29065 93.385414 611.2438 94.21354Z" fill-rule="nonzero"/><path fill="#fbbc04" d="M577.65 117.916664C577.65 121.23038 574.9637 123.916664 571.65 123.916664C568.3363 123.916664 565.65 121.23038 565.65 117.916664C565.65 114.60296 568.3363 111.916664 571.65 111.916664C574.9637 111.916664 577.65 114.60296 577.65 117.916664Z" fill-rule="nonzero"/><path fill="#1a1a1a" d="M588.0719 121.916664L585.79065 113.322914L586.9625 113.322914L588.275 118.96354Q588.47815 119.83854 588.6344 120.71354Q588.9625 119.33854 589.025 119.11979L590.65 113.322914L592.025 113.322914L593.2594 117.666664Q593.7125 119.291664 593.9313 120.71354Q594.0875 119.90104 594.35315 118.83854L595.6969 113.322914L596.8375 113.322914L594.4938 121.916664L593.3844 121.916664L591.5875 115.36979Q591.35315 114.55729 591.3063 114.36979Q591.1813 114.947914 591.0563 115.36979L589.2438 121.916664L588.0719 121.916664ZM597.0563 118.80729Q597.0563 117.072914 598.0094 116.24479Q598.8063 115.55729 599.9625 115.55729Q601.2438 115.55729 602.0563 116.40104Q602.8844 117.229164 602.8844 118.71354Q602.8844 119.916664 602.525 120.604164Q602.16565 121.291664 601.4625 121.68229Q600.775 122.05729 599.9625 122.05729Q598.66565 122.05729 597.85315 121.229164Q597.0563 120.385414 597.0563 118.80729ZM598.1344 118.80729Q598.1344 119.99479 598.65 120.58854Q599.1813 121.18229 599.9625 121.18229Q600.7594 121.18229 601.275 120.58854Q601.79065 119.99479 601.79065 118.77604Q601.79065 117.61979 601.2594 117.02604Q600.7438 116.43229 599.9625 116.43229Q599.1813 116.43229 598.65 117.02604Q598.1344 117.604164 598.1344 118.80729ZM604.4313 121.916664L604.4313 115.697914L605.3844 115.697914L605.3844 116.635414Q605.7438 115.979164 606.04065 115.77604Q606.35315 115.55729 606.72815 115.55729Q607.2594 115.55729 607.8063 115.885414L607.4469 116.86979Q607.0563 116.635414 606.6813 116.635414Q606.3219 116.635414 606.04065 116.854164Q605.775 117.05729 605.66565 117.43229Q605.47815 117.99479 605.47815 118.65104L605.47815 121.916664L604.4313 121.916664ZM608.4469 121.916664L608.4469 113.322914L609.5094 113.322914L609.5094 118.229164L611.9938 115.697914L613.3688 115.697914L610.97815 117.99479L613.60315 121.916664L612.3063 121.916664L610.2438 118.729164L609.5094 119.447914L609.5094 121.916664L608.4469 121.916664ZM621.6969 119.916664L622.79065 120.041664Q622.54065 120.99479 621.8375 121.52604Q621.1344 122.05729 620.0563 122.05729Q618.6969 122.05729 617.8844 121.21354Q617.0875 120.36979 617.0875 118.854164Q617.0875 117.291664 617.9 116.43229Q618.7125 115.55729 619.9938 115.55729Q621.2438 115.55729 622.025 116.40104Q622.8219 117.24479 622.8219 118.791664Q622.8219 118.885414 622.8219 119.072914L618.1813 119.072914Q618.2438 120.104164 618.7594 120.65104Q619.275 121.18229 620.0563 121.18229Q620.6344 121.18229 621.04065 120.885414Q621.4625 120.572914 621.6969 119.916664ZM618.2438 118.21354L621.7125 118.21354Q621.65 117.416664 621.3219 117.02604Q620.8063 116.416664 620.0094 116.416664Q619.275 116.416664 618.775 116.90104Q618.29065 117.385414 618.2438 118.21354ZM624.6969 121.916664L624.6969 116.510414L623.7594 116.510414L623.7594 115.697914L624.6969 115.697914L624.6969 115.02604Q624.6969 114.40104 624.8063 114.104164Q624.9625 113.68229 625.3375 113.43229Q625.72815 113.18229 626.41565 113.18229Q626.8688 113.18229 627.4 113.291664L627.2438 114.197914Q626.91565 114.15104 626.6188 114.15104Q626.1344 114.15104 625.9313 114.354164Q625.7438 114.55729 625.7438 115.11979L625.7438 115.697914L626.9625 115.697914L626.9625 116.510414L625.7438 116.510414L625.7438 121.916664L624.6969 121.916664ZM627.6969 121.916664L627.6969 116.510414L626.7594 116.510414L626.7594 115.697914L627.6969 115.697914L627.6969 115.02604Q627.6969 114.40104 627.8063 114.104164Q627.9625 113.68229 628.3375 113.43229Q628.72815 113.18229 629.41565 113.18229Q629.8688 113.18229 630.4 113.291664L630.2438 114.197914Q629.91565 114.15104 629.6188 114.15104Q629.1344 114.15104 628.9313 114.354164Q628.7438 114.55729 628.7438 115.11979L628.7438 115.697914L629.9625 115.697914L629.9625 116.510414L628.7438 116.510414L628.7438 121.916664L627.6969 121.916664ZM630.4469 114.541664L630.4469 113.322914L631.5094 113.322914L631.5094 114.541664L630.4469 114.541664ZM630.4469 121.916664L630.4469 115.697914L631.5094 115.697914L631.5094 121.916664L630.4469 121.916664ZM637.5094 119.635414L638.54065 119.77604Q638.3688 120.83854 637.66565 121.447914Q636.9625 122.05729 635.9469 122.05729Q634.66565 122.05729 633.8844 121.229164Q633.1188 120.385414 633.1188 118.822914Q633.1188 117.822914 633.4469 117.072914Q633.79065 116.30729 634.4625 115.93229Q635.15 115.55729 635.9625 115.55729Q636.9625 115.55729 637.60315 116.072914Q638.2594 116.572914 638.4469 117.52604L637.41565 117.68229Q637.275 117.05729 636.9 116.74479Q636.525 116.416664 635.9938 116.416664Q635.1969 116.416664 634.6969 116.99479Q634.1969 117.55729 634.1969 118.791664Q634.1969 120.05729 634.6813 120.61979Q635.16565 121.18229 635.9313 121.18229Q636.5563 121.18229 636.9625 120.80729Q637.3844 120.43229 637.5094 119.635414ZM639.4469 114.541664L639.4469 113.322914L640.5094 113.322914L640.5094 114.541664L639.4469 114.541664ZM639.4469 121.916664L639.4469 115.697914L640.5094 115.697914L640.5094 121.916664L639.4469 121.916664ZM646.6969 119.916664L647.79065 120.041664Q647.54065 120.99479 646.8375 121.52604Q646.1344 122.05729 645.0563 122.05729Q643.6969 122.05729 642.8844 121.21354Q642.0875 120.36979 642.0875 118.854164Q642.0875 117.291664 642.9 116.43229Q643.7125 115.55729 644.9938 115.55729Q646.2438 115.55729 647.025 116.40104Q647.8219 117.24479 647.8219 118.791664Q647.8219 118.885414 647.8219 119.072914L643.1813 119.072914Q643.2438 120.104164 643.7594 120.65104Q644.275 121.18229 645.0563 121.18229Q645.6344 121.18229 646.04065 120.885414Q646.4625 120.572914 646.6969 119.916664ZM643.2438 118.21354L646.7125 118.21354Q646.65 117.416664 646.3219 117.02604Q645.8063 116.416664 645.0094 116.416664Q644.275 116.416664 643.775 116.90104Q643.29065 117.385414 643.2438 118.21354ZM649.4469 121.916664L649.4469 115.697914L650.3844 115.697914L650.3844 116.572914Q651.0719 115.55729 652.3688 115.55729Q652.9313 115.55729 653.4 115.760414Q653.8844 115.96354 654.1188 116.291664Q654.35315 116.61979 654.4469 117.05729Q654.4938 117.354164 654.4938 118.08854L654.4938 121.916664L653.4469 121.916664L653.4469 118.135414Q653.4469 117.479164 653.3219 117.166664Q653.1969 116.854164 652.8844 116.666664Q652.5719 116.46354 652.15 116.46354Q651.47815 116.46354 650.97815 116.90104Q650.4938 117.322914 650.4938 118.510414L650.4938 121.916664L649.4469 121.916664ZM658.7438 120.979164L658.9 121.90104Q658.4469 121.99479 658.10315 121.99479Q657.525 121.99479 657.2125 121.822914Q656.9 121.635414 656.7594 121.33854Q656.6344 121.041664 656.6344 120.08854L656.6344 116.510414L655.8688 116.510414L655.8688 115.697914L656.6344 115.697914L656.6344 114.15104L657.6813 113.52604L657.6813 115.697914L658.7438 115.697914L658.7438 116.510414L657.6813 116.510414L657.6813 120.15104Q657.6813 120.604164 657.72815 120.729164Q657.79065 120.854164 657.91565 120.93229Q658.04065 121.010414 658.275 121.010414Q658.4625 121.010414 658.7438 120.979164Z" fill-rule="nonzero"/><path fill="#34a853" d="M577.65 141.91667C577.65 145.23038 574.9637 147.91667 571.65 147.91667C568.3363 147.91667 565.65 145.23038 565.65 141.91667C565.65 138.60295 568.3363 135.91667 571.65 135.91667C574.9637 135.91667 577.65 138.60295 577.65 141.91667Z" fill-rule="nonzero"/><path fill="#1a1a1a" d="M588.7594 145.91667L588.7594 138.33855L585.9313 138.33855L585.9313 137.32292L592.7438 137.32292L592.7438 138.33855L589.9 138.33855L589.9 145.91667L588.7594 145.91667ZM593.4469 145.91667L593.4469 137.32292L594.4938 137.32292L594.4938 140.40105Q595.22815 139.5573 596.35315 139.5573Q597.0563 139.5573 597.5563 139.83855Q598.0719 140.10417 598.29065 140.58855Q598.5094 141.0573 598.5094 141.97917L598.5094 145.91667L597.4625 145.91667L597.4625 141.97917Q597.4625 141.1823 597.1188 140.82292Q596.775 140.46355 596.15 140.46355Q595.6813 140.46355 595.2594 140.71355Q594.85315 140.94792 594.66565 141.3698Q594.4938 141.77605 594.4938 142.51042L594.4938 145.91667L593.4469 145.91667ZM600.4313 145.91667L600.4313 139.69792L601.3844 139.69792L601.3844 140.63542Q601.7438 139.97917 602.04065 139.77605Q602.35315 139.5573 602.72815 139.5573Q603.2594 139.5573 603.8063 139.88542L603.4469 140.8698Q603.0563 140.63542 602.6813 140.63542Q602.3219 140.63542 602.04065 140.85417Q601.775 141.0573 601.66565 141.4323Q601.47815 141.9948 601.47815 142.65105L601.47815 145.91667L600.4313 145.91667ZM608.525 145.91667L608.525 144.9948Q607.79065 146.0573 606.54065 146.0573Q605.9938 146.0573 605.5094 145.85417Q605.04065 145.63542 604.8063 145.32292Q604.5719 144.9948 604.47815 144.52605Q604.41565 144.22917 604.41565 143.54167L604.41565 139.69792L605.47815 139.69792L605.47815 143.15105Q605.47815 143.96355 605.54065 144.26042Q605.6344 144.66667 605.9469 144.91667Q606.275 145.15105 606.7594 145.15105Q607.22815 145.15105 607.6344 144.91667Q608.0563 144.66667 608.22815 144.2448Q608.41565 143.82292 608.41565 143.02605L608.41565 139.69792L609.4625 139.69792L609.4625 145.91667L608.525 145.91667ZM611.025 144.0573L612.0563 143.90105Q612.15 144.52605 612.54065 144.85417Q612.9469 145.1823 613.6813 145.1823Q614.4 145.1823 614.7438 144.90105Q615.10315 144.60417 615.10315 144.19792Q615.10315 143.83855 614.79065 143.63542Q614.5719 143.4948 613.7125 143.27605Q612.5563 142.97917 612.10315 142.77605Q611.66565 142.5573 611.4313 142.1823Q611.1969 141.8073 611.1969 141.33855Q611.1969 140.9323 611.3844 140.57292Q611.5719 140.21355 611.9 139.97917Q612.15 139.8073 612.5719 139.6823Q612.9938 139.5573 613.4938 139.5573Q614.2125 139.5573 614.7594 139.77605Q615.3219 139.97917 615.5875 140.33855Q615.85315 140.69792 615.9469 141.29167L614.91565 141.4323Q614.85315 140.96355 614.5094 140.69792Q614.1813 140.41667 613.5563 140.41667Q612.8375 140.41667 612.525 140.66667Q612.2125 140.90105 612.2125 141.22917Q612.2125 141.4323 612.3375 141.58855Q612.47815 141.76042 612.7438 141.8698Q612.9 141.9323 613.6813 142.13542Q614.8063 142.44792 615.2438 142.63542Q615.6813 142.82292 615.9313 143.1823Q616.1813 143.54167 616.1813 144.08855Q616.1813 144.6198 615.8688 145.08855Q615.5719 145.54167 614.9938 145.8073Q614.41565 146.0573 613.6813 146.0573Q612.4625 146.0573 611.8219 145.5573Q611.1969 145.04167 611.025 144.0573ZM620.7438 144.97917L620.9 145.90105Q620.4469 145.9948 620.10315 145.9948Q619.525 145.9948 619.2125 145.82292Q618.9 145.63542 618.7594 145.33855Q618.6344 145.04167 618.6344 144.08855L618.6344 140.51042L617.8688 140.51042L617.8688 139.69792L618.6344 139.69792L618.6344 138.15105L619.6813 137.52605L619.6813 139.69792L620.7438 139.69792L620.7438 140.51042L619.6813 140.51042L619.6813 144.15105Q619.6813 144.60417 619.72815 144.72917Q619.79065 144.85417 619.91565 144.9323Q620.04065 145.01042 620.275 145.01042Q620.4625 145.01042 620.7438 144.97917Z" fill-rule="nonzero"/><path fill="#757575" d="M24.896875 42.35L24.896875 28.037498L30.2875 28.037498Q31.725 28.037498 32.475 28.162498Q33.521873 28.349998 34.240623 28.849998Q34.959373 29.334373 35.38125 30.224998Q35.81875 31.115623 35.81875 32.178123Q35.81875 34.00625 34.646873 35.271873Q33.490623 36.521873 30.459375 36.521873L26.7875 36.521873L26.7875 42.35L24.896875 42.35ZM26.7875 34.834373L30.490625 34.834373Q32.31875 34.834373 33.084373 34.1625Q33.865623 33.475 33.865623 32.225Q33.865623 31.334373 33.4125 30.693748Q32.959373 30.053123 32.225 29.849998Q31.740625 29.724998 30.44375 29.724998L26.7875 29.724998L26.7875 34.834373ZM44.771873 39.00625L46.584373 39.240623Q46.1625 40.81875 44.990623 41.709373Q43.834373 42.584373 42.021873 42.584373Q39.75625 42.584373 38.4125 41.19375Q37.084373 39.7875 37.084373 37.25625Q37.084373 34.63125 38.428123 33.19375Q39.771873 31.740623 41.928123 31.740623Q44.00625 31.740623 45.31875 33.1625Q46.646873 34.56875 46.646873 37.146873Q46.646873 37.303123 46.63125 37.615623L38.896873 37.615623Q38.990623 39.31875 39.865623 40.225Q40.740623 41.13125 42.0375 41.13125Q43.00625 41.13125 43.678123 40.63125Q44.365623 40.115623 44.771873 39.00625ZM38.990623 36.1625L44.7875 36.1625Q44.678123 34.865623 44.13125 34.209373Q43.2875 33.19375 41.94375 33.19375Q40.740623 33.19375 39.9125 34.00625Q39.084373 34.803123 38.990623 36.1625ZM48.646873 42.35L48.646873 31.974998L50.225 31.974998L50.225 33.553123Q50.834373 32.44375 51.35 32.1Q51.865623 31.740623 52.475 31.740623Q53.365623 31.740623 54.2875 32.303123L53.678123 33.94375Q53.0375 33.553123 52.396873 33.553123Q51.81875 33.553123 51.35 33.9125Q50.896873 34.25625 50.69375 34.865623Q50.4125 35.803123 50.4125 36.9125L50.4125 42.35L48.646873 42.35ZM56.084373 42.35L56.084373 33.35L54.5375 33.35L54.5375 31.974998L56.084373 31.974998L56.084373 30.881248Q56.084373 29.834373 56.271873 29.318748Q56.521873 28.631248 57.1625 28.209373Q57.803123 27.787498 58.959373 27.787498Q59.709373 27.787498 60.6 27.959373L60.334373 29.490623Q59.7875 29.396873 59.303123 29.396873Q58.50625 29.396873 58.1625 29.740623Q57.834373 30.084373 57.834373 31.021873L57.834373 31.974998L59.865623 31.974998L59.865623 33.35L57.834373 33.35L57.834373 42.35L56.084373 42.35ZM61.021873 37.1625Q61.021873 34.2875 62.615623 32.896873Q63.959373 31.740623 65.88125 31.740623Q68.02187 31.740623 69.36562 33.146873Q70.725 34.553123 70.725 37.021873Q70.725 39.021873 70.13125 40.178123Q69.5375 41.31875 68.38125 41.959373Q67.24062 42.584373 65.88125 42.584373Q63.69375 42.584373 62.35 41.19375Q61.021873 39.7875 61.021873 37.1625ZM62.81875 37.1625Q62.81875 39.1625 63.678123 40.146873Q64.55312 41.13125 65.88125 41.13125Q67.19375 41.13125 68.05312 40.146873Q68.92812 39.146873 68.92812 37.1Q68.92812 35.178123 68.05312 34.19375Q67.17812 33.19375 65.88125 33.19375Q64.55312 33.19375 63.678123 34.19375Q62.81875 35.178123 62.81875 37.1625ZM72.64687 42.35L72.64687 31.974998L74.225 31.974998L74.225 33.553123Q74.83437 32.44375 75.35 32.1Q75.86562 31.740623 76.475 31.740623Q77.36562 31.740623 78.2875 32.303123L77.67812 33.94375Q77.0375 33.553123 76.39687 33.553123Q75.81875 33.553123 75.35 33.9125Q74.89687 34.25625 74.69375 34.865623Q74.4125 35.803123 74.4125 36.9125L74.4125 42.35L72.64687 42.35ZM79.6625 42.35L79.6625 31.974998L81.24062 31.974998L81.24062 33.428123Q81.725 32.678123 82.5375 32.209373Q83.35 31.740623 84.38125 31.740623Q85.5375 31.740623 86.27187 32.225Q87.00625 32.69375 87.31875 33.553123Q88.55312 31.740623 90.52187 31.740623Q92.06875 31.740623 92.89687 32.6Q93.725 33.459373 93.725 35.225L93.725 42.35L91.975 42.35L91.975 35.81875Q91.975 34.75625 91.80312 34.303123Q91.63125 33.834373 91.17812 33.553123Q90.74062 33.271873 90.13125 33.271873Q89.0375 33.271873 88.30312 34.00625Q87.58437 34.725 87.58437 36.31875L87.58437 42.35L85.83437 42.35L85.83437 35.615623Q85.83437 34.44375 85.39687 33.865623Q84.975 33.271873 83.99062 33.271873Q83.25625 33.271873 82.61562 33.6625Q81.99062 34.053123 81.70937 34.803123Q81.42812 35.553123 81.42812 36.975L81.42812 42.35L79.6625 42.35ZM102.44375 41.06875Q101.45937 41.896873 100.55312 42.240623Q99.64687 42.584373 98.61562 42.584373Q96.9125 42.584373 95.99062 41.75625Q95.06875 40.9125 95.06875 39.615623Q95.06875 38.85 95.4125 38.225Q95.77187 37.6 96.33437 37.225Q96.89687 36.834373 97.6 36.63125Q98.11562 36.50625 99.1625 36.365623Q101.2875 36.115623 102.2875 35.771873Q102.30312 35.4125 102.30312 35.303123Q102.30312 34.240623 101.80312 33.7875Q101.13125 33.19375 99.80312 33.19375Q98.55312 33.19375 97.95937 33.63125Q97.38125 34.06875 97.1 35.178123L95.38125 34.94375Q95.61562 33.834373 96.14687 33.1625Q96.67812 32.475 97.69375 32.115623Q98.70937 31.740623 100.05312 31.740623Q101.38125 31.740623 102.20937 32.053123Q103.0375 32.365623 103.42812 32.85Q103.81875 33.31875 103.975 34.0375Q104.06875 34.490623 104.06875 35.6625L104.06875 38.00625Q104.06875 40.459373 104.17812 41.115623Q104.2875 41.75625 104.63125 42.35L102.7875 42.35Q102.52187 41.803123 102.44375 41.06875ZM102.2875 37.146873Q101.33437 37.5375 99.4125 37.803123Q98.33437 37.959373 97.88125 38.1625Q97.44375 38.35 97.19375 38.725Q96.94375 39.1 96.94375 39.56875Q96.94375 40.271873 97.475 40.740623Q98.00625 41.209373 99.0375 41.209373Q100.05312 41.209373 100.83437 40.771873Q101.63125 40.31875 102.00625 39.553123Q102.2875 38.94375 102.2875 37.7875L102.2875 37.146873ZM106.6625 42.35L106.6625 31.974998L108.25625 31.974998L108.25625 33.459373Q109.39687 31.740623 111.55312 31.740623Q112.49062 31.740623 113.27187 32.084373Q114.06875 32.4125 114.45937 32.959373Q114.85 33.50625 115.00625 34.25625Q115.1 34.75625 115.1 35.975L115.1 42.35L113.33437 42.35L113.33437 36.0375Q113.33437 34.959373 113.13125 34.428123Q112.92812 33.896873 112.39687 33.584373Q111.88125 33.271873 111.17812 33.271873Q110.05312 33.271873 109.24062 33.990623Q108.42812 34.69375 108.42812 36.678123L108.42812 42.35L106.6625 42.35ZM123.44375 38.553123L125.1625 38.771873Q124.88125 40.56875 123.70937 41.584373Q122.55312 42.584373 120.85 42.584373Q118.725 42.584373 117.42812 41.19375Q116.13125 39.803123 116.13125 37.209373Q116.13125 35.521873 116.67812 34.271873Q117.24062 33.00625 118.38125 32.38125Q119.52187 31.740623 120.86562 31.740623Q122.55312 31.740623 123.61562 32.6Q124.69375 33.459373 125.00625 35.021873L123.2875 35.2875Q123.05312 34.240623 122.42812 33.725Q121.80312 33.19375 120.92812 33.19375Q119.6 33.19375 118.77187 34.146873Q117.94375 35.1 117.94375 37.146873Q117.94375 39.240623 118.74062 40.19375Q119.5375 41.13125 120.83437 41.13125Q121.86562 41.13125 122.55312 40.50625Q123.25625 39.865623 123.44375 38.553123ZM133.77188 39.00625L135.58438 39.240623Q135.1625 40.81875 133.99063 41.709373Q132.83438 42.584373 131.02188 42.584373Q128.75626 42.584373 127.4125 41.19375Q126.08437 39.7875 126.08437 37.25625Q126.08437 34.63125 127.42812 33.19375Q128.77188 31.740623 130.92813 31.740623Q133.00626 31.740623 134.31876 33.1625Q135.64688 34.56875 135.64688 37.146873Q135.64688 37.303123 135.63126 37.615623L127.89687 37.615623Q127.99062 39.31875 128.86563 40.225Q129.74063 41.13125 131.0375 41.13125Q132.00626 41.13125 132.67813 40.63125Q133.36563 40.115623 133.77188 39.00625ZM127.99062 36.1625L133.7875 36.1625Q133.67813 34.865623 133.13126 34.209373Q132.2875 33.19375 130.94376 33.19375Q129.74063 33.19375 128.9125 34.00625Q128.08438 34.803123 127.99062 36.1625ZM143.02188 37.1625Q143.02188 34.2875 144.61563 32.896873Q145.95938 31.740623 147.88126 31.740623Q150.02188 31.740623 151.36563 33.146873Q152.725 34.553123 152.725 37.021873Q152.725 39.021873 152.13126 40.178123Q151.5375 41.31875 150.38126 41.959373Q149.24063 42.584373 147.88126 42.584373Q145.69376 42.584373 144.35 41.19375Q143.02188 39.7875 143.02188 37.1625ZM144.81876 37.1625Q144.81876 39.1625 145.67813 40.146873Q146.55313 41.13125 147.88126 41.13125Q149.19376 41.13125 150.05313 40.146873Q150.92813 39.146873 150.92813 37.1Q150.92813 35.178123 150.05313 34.19375Q149.17813 33.19375 147.88126 33.19375Q146.55313 33.19375 145.67813 34.19375Q144.81876 35.178123 144.81876 37.1625ZM155.08438 42.35L155.08438 33.35L153.5375 33.35L153.5375 31.974998L155.08438 31.974998L155.08438 30.881248Q155.08438 29.834373 155.27188 29.318748Q155.52188 28.631248 156.1625 28.209373Q156.80313 27.787498 157.95938 27.787498Q158.70938 27.787498 159.6 27.959373L159.33438 29.490623Q158.7875 29.396873 158.30313 29.396873Q157.50626 29.396873 157.1625 29.740623Q156.83438 30.084373 156.83438 31.021873L156.83438 31.974998L158.86563 31.974998L158.86563 33.35L156.83438 33.35L156.83438 42.35L155.08438 42.35ZM166.25626 37.75625L168.0375 37.6Q168.1625 38.6625 168.61563 39.35Q169.08438 40.0375 170.05313 40.475Q171.0375 40.896873 172.27188 40.896873Q173.35 40.896873 174.17813 40.584373Q175.00626 40.25625 175.4125 39.69375Q175.81876 39.13125 175.81876 38.459373Q175.81876 37.7875 175.42813 37.2875Q175.0375 36.7875 174.14688 36.44375Q173.56876 36.209373 171.58438 35.740623Q169.61563 35.271873 168.83438 34.85Q167.80313 34.31875 167.30313 33.521873Q166.80313 32.725 166.80313 31.740623Q166.80313 30.646873 167.4125 29.709373Q168.02188 28.771873 169.20938 28.287498Q170.39688 27.787498 171.83438 27.787498Q173.42813 27.787498 174.64688 28.303123Q175.86563 28.818748 176.50626 29.818748Q177.1625 30.803123 177.20938 32.06875L175.39688 32.209373Q175.25626 30.849998 174.4125 30.162498Q173.56876 29.459373 171.9125 29.459373Q170.19376 29.459373 169.39688 30.099998Q168.61563 30.724998 168.61563 31.599998Q168.61563 32.38125 169.17813 32.88125Q169.725 33.38125 172.0375 33.896873Q174.35 34.4125 175.20938 34.803123Q176.45938 35.38125 177.05313 36.271873Q177.64688 37.146873 177.64688 38.303123Q177.64688 39.44375 176.99063 40.459373Q176.33438 41.475 175.1 42.0375Q173.88126 42.6 172.35 42.6Q170.4125 42.6 169.1 42.0375Q167.7875 41.459373 167.0375 40.31875Q166.2875 39.178123 166.25626 37.75625ZM186.44376 38.553123L188.1625 38.771873Q187.88126 40.56875 186.70938 41.584373Q185.55313 42.584373 183.85 42.584373Q181.725 42.584373 180.42813 41.19375Q179.13126 39.803123 179.13126 37.209373Q179.13126 35.521873 179.67813 34.271873Q180.24063 33.00625 181.38126 32.38125Q182.52188 31.740623 183.86563 31.740623Q185.55313 31.740623 186.61563 32.6Q187.69376 33.459373 188.00626 35.021873L186.2875 35.2875Q186.05313 34.240623 185.42813 33.725Q184.80313 33.19375 183.92813 33.19375Q182.6 33.19375 181.77188 34.146873Q180.94376 35.1 180.94376 37.146873Q180.94376 39.240623 181.74063 40.19375Q182.5375 41.13125 183.83438 41.13125Q184.86563 41.13125 185.55313 40.50625Q186.25626 39.865623 186.44376 38.553123ZM196.44376 41.06875Q195.45938 41.896873 194.55313 42.240623Q193.64688 42.584373 192.61563 42.584373Q190.9125 42.584373 189.99063 41.75625Q189.06876 40.9125 189.06876 39.615623Q189.06876 38.85 189.4125 38.225Q189.77188 37.6 190.33438 37.225Q190.89688 36.834373 191.6 36.63125Q192.11563 36.50625 193.1625 36.365623Q195.2875 36.115623 196.2875 35.771873Q196.30313 35.4125 196.30313 35.303123Q196.30313 34.240623 195.80313 33.7875Q195.13126 33.19375 193.80313 33.19375Q192.55313 33.19375 191.95938 33.63125Q191.38126 34.06875 191.1 35.178123L189.38126 34.94375Q189.61563 33.834373 190.14688 33.1625Q190.67813 32.475 191.69376 32.115623Q192.70938 31.740623 194.05313 31.740623Q195.38126 31.740623 196.20938 32.053123Q197.0375 32.365623 197.42813 32.85Q197.81876 33.31875 197.975 34.0375Q198.06876 34.490623 198.06876 35.6625L198.06876 38.00625Q198.06876 40.459373 198.17813 41.115623Q198.2875 41.75625 198.63126 42.35L196.7875 42.35Q196.52188 41.803123 196.44376 41.06875ZM196.2875 37.146873Q195.33438 37.5375 193.4125 37.803123Q192.33438 37.959373 191.88126 38.1625Q191.44376 38.35 191.19376 38.725Q190.94376 39.1 190.94376 39.56875Q190.94376 40.271873 191.475 40.740623Q192.00626 41.209373 193.0375 41.209373Q194.05313 41.209373 194.83438 40.771873Q195.63126 40.31875 196.00626 39.553123Q196.2875 38.94375 196.2875 37.7875L196.2875 37.146873ZM200.6625 42.35L200.6625 31.974998L202.25626 31.974998L202.25626 33.459373Q203.39688 31.740623 205.55313 31.740623Q206.49063 31.740623 207.27188 32.084373Q208.06876 32.4125 208.45938 32.959373Q208.85 33.50625 209.00626 34.25625Q209.1 34.75625 209.1 35.975L209.1 42.35L207.33438 42.35L207.33438 36.0375Q207.33438 34.959373 207.13126 34.428123Q206.92813 33.896873 206.39688 33.584373Q205.88126 33.271873 205.17813 33.271873Q204.05313 33.271873 203.24063 33.990623Q202.42813 34.69375 202.42813 36.678123L202.42813 42.35L200.6625 42.35ZM217.20938 42.35L217.20938 28.037498L219.11563 28.037498L219.11563 42.35L217.20938 42.35ZM222.6625 42.35L222.6625 31.974998L224.24063 31.974998L224.24063 33.428123Q224.725 32.678123 225.5375 32.209373Q226.35 31.740623 227.38126 31.740623Q228.5375 31.740623 229.27188 32.225Q230.00626 32.69375 230.31876 33.553123Q231.55313 31.740623 233.52188 31.740623Q235.06876 31.740623 235.89688 32.6Q236.725 33.459373 236.725 35.225L236.725 42.35L234.975 42.35L234.975 35.81875Q234.975 34.75625 234.80313 34.303123Q234.63126 33.834373 234.17813 33.553123Q233.74063 33.271873 233.13126 33.271873Q232.0375 33.271873 231.30313 34.00625Q230.58438 34.725 230.58438 36.31875L230.58438 42.35L228.83438 42.35L228.83438 35.615623Q228.83438 34.44375 228.39688 33.865623Q227.975 33.271873 226.99063 33.271873Q226.25626 33.271873 225.61563 33.6625Q224.99063 34.053123 224.70938 34.803123Q224.42813 35.553123 224.42813 36.975L224.42813 42.35L222.6625 42.35ZM238.6625 46.31875L238.6625 31.974998L240.27188 31.974998L240.27188 33.31875Q240.83438 32.5375 241.5375 32.146873Q242.25626 31.740623 243.27188 31.740623Q244.6 31.740623 245.61563 32.428123Q246.63126 33.115623 247.14688 34.365623Q247.67813 35.6 247.67813 37.084373Q247.67813 38.678123 247.1 39.959373Q246.5375 41.225 245.44376 41.9125Q244.35 42.584373 243.14688 42.584373Q242.27188 42.584373 241.56876 42.209373Q240.88126 41.834373 240.42813 41.271873L240.42813 46.31875L238.6625 46.31875ZM240.25626 37.225Q240.25626 39.225 241.06876 40.178123Q241.88126 41.13125 243.0375 41.13125Q244.20938 41.13125 245.0375 40.146873Q245.88126 39.1625 245.88126 37.06875Q245.88126 35.1 245.05313 34.115623Q244.24063 33.115623 243.11563 33.115623Q241.99063 33.115623 241.11563 34.178123Q240.25626 35.225 240.25626 37.225ZM249.63126 42.35L249.63126 28.037498L251.38126 28.037498L251.38126 42.35L249.63126 42.35ZM260.77188 39.00625L262.58438 39.240623Q262.1625 40.81875 260.99063 41.709373Q259.83438 42.584373 258.02188 42.584373Q255.75626 42.584373 254.4125 41.19375Q253.08438 39.7875 253.08438 37.25625Q253.08438 34.63125 254.42813 33.19375Q255.77188 31.740623 257.92813 31.740623Q260.00626 31.740623 261.31876 33.1625Q262.64688 34.56875 262.64688 37.146873Q262.64688 37.303123 262.63126 37.615623L254.89688 37.615623Q254.99063 39.31875 255.86563 40.225Q256.74063 41.13125 258.0375 41.13125Q259.00626 41.13125 259.67813 40.63125Q260.36563 40.115623 260.77188 39.00625ZM254.99063 36.1625L260.7875 36.1625Q260.67813 34.865623 260.13126 34.209373Q259.2875 33.19375 257.94376 33.19375Q256.74063 33.19375 255.9125 34.00625Q255.08438 34.803123 254.99063 36.1625ZM264.6625 42.35L264.6625 31.974998L266.24063 31.974998L266.24063 33.428123Q266.725 32.678123 267.5375 32.209373Q268.35 31.740623 269.38126 31.740623Q270.5375 31.740623 271.27188 32.225Q272.00626 32.69375 272.31876 33.553123Q273.55313 31.740623 275.52188 31.740623Q277.06876 31.740623 277.89688 32.6Q278.725 33.459373 278.725 35.225L278.725 42.35L276.975 42.35L276.975 35.81875Q276.975 34.75625 276.80313 34.303123Q276.63126 33.834373 276.17813 33.553123Q275.74063 33.271873 275.13126 33.271873Q274.0375 33.271873 273.30313 34.00625Q272.58438 34.725 272.58438 36.31875L272.58438 42.35L270.83438 42.35L270.83438 35.615623Q270.83438 34.44375 270.39688 33.865623Q269.975 33.271873 268.99063 33.271873Q268.25626 33.271873 267.61563 33.6625Q266.99063 34.053123 266.70938 34.803123Q266.42813 35.553123 266.42813 36.975L266.42813 42.35L264.6625 42.35ZM287.77188 39.00625L289.58438 39.240623Q289.1625 40.81875 287.99063 41.709373Q286.83438 42.584373 285.02188 42.584373Q282.75626 42.584373 281.4125 41.19375Q280.08438 39.7875 280.08438 37.25625Q280.08438 34.63125 281.42813 33.19375Q282.77188 31.740623 284.92813 31.740623Q287.00626 31.740623 288.31876 33.1625Q289.64688 34.56875 289.64688 37.146873Q289.64688 37.303123 289.63126 37.615623L281.89688 37.615623Q281.99063 39.31875 282.86563 40.225Q283.74063 41.13125 285.0375 41.13125Q286.00626 41.13125 286.67813 40.63125Q287.36563 40.115623 287.77188 39.00625ZM281.99063 36.1625L287.7875 36.1625Q287.67813 34.865623 287.13126 34.209373Q286.2875 33.19375 284.94376 33.19375Q283.74063 33.19375 282.9125 34.00625Q282.08438 34.803123 281.99063 36.1625ZM291.6625 42.35L291.6625 31.974998L293.25626 31.974998L293.25626 33.459373Q294.39688 31.740623 296.55313 31.740623Q297.49063 31.740623 298.27188 32.084373Q299.06876 32.4125 299.45938 32.959373Q299.85 33.50625 300.00626 34.25625Q300.1 34.75625 300.1 35.975L300.1 42.35L298.33438 42.35L298.33438 36.0375Q298.33438 34.959373 298.13126 34.428123Q297.92813 33.896873 297.39688 33.584373Q296.88126 33.271873 296.17813 33.271873Q295.05313 33.271873 294.24063 33.990623Q293.42813 34.69375 293.42813 36.678123L293.42813 42.35L291.6625 42.35ZM305.50626 40.771873L305.75626 42.334373Q305.02188 42.490623 304.42813 42.490623Q303.475 42.490623 302.94376 42.19375Q302.42813 41.88125 302.20938 41.396873Q301.99063 40.896873 301.99063 39.31875L301.99063 33.35L300.70938 33.35L300.70938 31.974998L301.99063 31.974998L301.99063 29.412498L303.74063 28.349998L303.74063 31.974998L305.50626 31.974998L305.50626 33.35L303.74063 33.35L303.74063 39.4125Q303.74063 40.1625 303.83438 40.38125Q303.92813 40.584373 304.13126 40.725Q304.35 40.85 304.74063 40.85Q305.02188 40.85 305.50626 40.771873ZM314.44376 41.06875Q313.45938 41.896873 312.55313 42.240623Q311.64688 42.584373 310.61563 42.584373Q308.9125 42.584373 307.99063 41.75625Q307.06876 40.9125 307.06876 39.615623Q307.06876 38.85 307.4125 38.225Q307.77188 37.6 308.33438 37.225Q308.89688 36.834373 309.6 36.63125Q310.11563 36.50625 311.1625 36.365623Q313.2875 36.115623 314.2875 35.771873Q314.30313 35.4125 314.30313 35.303123Q314.30313 34.240623 313.80313 33.7875Q313.13126 33.19375 311.80313 33.19375Q310.55313 33.19375 309.95938 33.63125Q309.38126 34.06875 309.1 35.178123L307.38126 34.94375Q307.61563 33.834373 308.14688 33.1625Q308.67813 32.475 309.69376 32.115623Q310.70938 31.740623 312.05313 31.740623Q313.38126 31.740623 314.20938 32.053123Q315.0375 32.365623 315.42813 32.85Q315.81876 33.31875 315.975 34.0375Q316.06876 34.490623 316.06876 35.6625L316.06876 38.00625Q316.06876 40.459373 316.17813 41.115623Q316.2875 41.75625 316.63126 42.35L314.7875 42.35Q314.52188 41.803123 314.44376 41.06875ZM314.2875 37.146873Q313.33438 37.5375 311.4125 37.803123Q310.33438 37.959373 309.88126 38.1625Q309.44376 38.35 309.19376 38.725Q308.94376 39.1 308.94376 39.56875Q308.94376 40.271873 309.475 40.740623Q310.00626 41.209373 311.0375 41.209373Q312.05313 41.209373 312.83438 40.771873Q313.63126 40.31875 314.00626 39.553123Q314.2875 38.94375 314.2875 37.7875L314.2875 37.146873ZM322.50626 40.771873L322.75626 42.334373Q322.02188 42.490623 321.42813 42.490623Q320.475 42.490623 319.94376 42.19375Q319.42813 41.88125 319.20938 41.396873Q318.99063 40.896873 318.99063 39.31875L318.99063 33.35L317.70938 33.35L317.70938 31.974998L318.99063 31.974998L318.99063 29.412498L320.74063 28.349998L320.74063 31.974998L322.50626 31.974998L322.50626 33.35L320.74063 33.35L320.74063 39.4125Q320.74063 40.1625 320.83438 40.38125Q320.92813 40.584373 321.13126 40.725Q321.35 40.85 321.74063 40.85Q322.02188 40.85 322.50626 40.771873ZM324.67813 30.053123L324.67813 28.037498L326.44376 28.037498L326.44376 30.053123L324.67813 30.053123ZM324.67813 42.35L324.67813 31.974998L326.44376 31.974998L326.44376 42.35L324.67813 42.35ZM328.02188 37.1625Q328.02188 34.2875 329.61563 32.896873Q330.95938 31.740623 332.88126 31.740623Q335.02188 31.740623 336.36563 33.146873Q337.725 34.553123 337.725 37.021873Q337.725 39.021873 337.13126 40.178123Q336.5375 41.31875 335.38126 41.959373Q334.24063 42.584373 332.88126 42.584373Q330.69376 42.584373 329.35 41.19375Q328.02188 39.7875 328.02188 37.1625ZM329.81876 37.1625Q329.81876 39.1625 330.67813 40.146873Q331.55313 41.13125 332.88126 41.13125Q334.19376 41.13125 335.05313 40.146873Q335.92813 39.146873 335.92813 37.1Q335.92813 35.178123 335.05313 34.19375Q334.17813 33.19375 332.88126 33.19375Q331.55313 33.19375 330.67813 34.19375Q329.81876 35.178123 329.81876 37.1625ZM339.6625 42.35L339.6625 31.974998L341.25626 31.974998L341.25626 33.459373Q342.39688 31.740623 344.55313 31.740623Q345.49063 31.740623 346.27188 32.084373Q347.06876 32.4125 347.45938 32.959373Q347.85 33.50625 348.00626 34.25625Q348.1 34.75625 348.1 35.975L348.1 42.35L346.33438 42.35L346.33438 36.0375Q346.33438 34.959373 346.13126 34.428123Q345.92813 33.896873 345.39688 33.584373Q344.88126 33.271873 344.17813 33.271873Q343.05313 33.271873 342.24063 33.990623Q341.42813 34.69375 341.42813 36.678123L341.42813 42.35L339.6625 42.35ZM348.95938 39.25625L350.70938 38.975Q350.85 40.021873 351.52188 40.584373Q352.19376 41.13125 353.39688 41.13125Q354.6 41.13125 355.17813 40.646873Q355.77188 40.146873 355.77188 39.490623Q355.77188 38.896873 355.25626 38.553123Q354.89688 38.31875 353.45938 37.959373Q351.52188 37.459373 350.77188 37.115623Q350.0375 36.75625 349.64688 36.13125Q349.25626 35.490623 349.25626 34.725Q349.25626 34.0375 349.56876 33.44375Q349.89688 32.85 350.44376 32.459373Q350.85 32.1625 351.55313 31.959373Q352.27188 31.740623 353.08438 31.740623Q354.30313 31.740623 355.225 32.1Q356.14688 32.44375 356.58438 33.053123Q357.02188 33.646873 357.19376 34.646873L355.475 34.88125Q355.35 34.084373 354.7875 33.646873Q354.225 33.19375 353.20938 33.19375Q351.99063 33.19375 351.475 33.6Q350.95938 33.990623 350.95938 34.521873Q350.95938 34.865623 351.17813 35.146873Q351.38126 35.428123 351.85 35.615623Q352.11563 35.709373 353.39688 36.053123Q355.27188 36.553123 356.00626 36.88125Q356.74063 37.19375 357.1625 37.803123Q357.58438 38.396873 357.58438 39.303123Q357.58438 40.178123 357.06876 40.959373Q356.55313 41.740623 355.58438 42.1625Q354.61563 42.584373 353.39688 42.584373Q351.38126 42.584373 350.31876 41.740623Q349.25626 40.896873 348.95938 39.25625ZM365.02188 37.1625Q365.02188 34.2875 366.61563 32.896873Q367.95938 31.740623 369.88126 31.740623Q372.02188 31.740623 373.36563 33.146873Q374.725 34.553123 374.725 37.021873Q374.725 39.021873 374.13126 40.178123Q373.5375 41.31875 372.38126 41.959373Q371.24063 42.584373 369.88126 42.584373Q367.69376 42.584373 366.35 41.19375Q365.02188 39.7875 365.02188 37.1625ZM366.81876 37.1625Q366.81876 39.1625 367.67813 40.146873Q368.55313 41.13125 369.88126 41.13125Q371.19376 41.13125 372.05313 40.146873Q372.92813 39.146873 372.92813 37.1Q372.92813 35.178123 372.05313 34.19375Q371.17813 33.19375 369.88126 33.19375Q368.55313 33.19375 367.67813 34.19375Q366.81876 35.178123 366.81876 37.1625ZM376.6625 42.35L376.6625 31.974998L378.25626 31.974998L378.25626 33.459373Q379.39688 31.740623 381.55313 31.740623Q382.49063 31.740623 383.27188 32.084373Q384.06876 32.4125 384.45938 32.959373Q384.85 33.50625 385.00626 34.25625Q385.1 34.75625 385.1 35.975L385.1 42.35L383.33438 42.35L383.33438 36.0375Q383.33438 34.959373 383.13126 34.428123Q382.92813 33.896873 382.39688 33.584373Q381.88126 33.271873 381.17813 33.271873Q380.05313 33.271873 379.24063 33.990623Q378.42813 34.69375 378.42813 36.678123L378.42813 42.35L376.6625 42.35ZM392.89688 42.35L392.89688 28.037498L397.81876 28.037498Q399.49063 28.037498 400.38126 28.240623Q401.6 28.521873 402.475 29.256248Q403.6 30.224998 404.1625 31.709373Q404.725 33.19375 404.725 35.115623Q404.725 36.740623 404.35 38.00625Q403.975 39.25625 403.36563 40.084373Q402.77188 40.9125 402.06876 41.396873Q401.36563 41.865623 400.36563 42.115623Q399.36563 42.35 398.05313 42.35L392.89688 42.35ZM394.7875 40.6625L397.85 40.6625Q399.25626 40.6625 400.05313 40.396873Q400.86563 40.13125 401.35 39.646873Q402.02188 38.975 402.39688 37.85Q402.77188 36.709373 402.77188 35.084373Q402.77188 32.834373 402.0375 31.631248Q401.30313 30.428123 400.24063 30.021873Q399.49063 29.724998 397.80313 29.724998L394.7875 29.724998L394.7875 40.6625ZM406.67813 30.053123L406.67813 28.037498L408.44376 28.037498L408.44376 30.053123L406.67813 30.053123ZM406.67813 42.35L406.67813 31.974998L408.44376 31.974998L408.44376 42.35L406.67813 42.35ZM411.08438 42.35L411.08438 33.35L409.5375 33.35L409.5375 31.974998L411.08438 31.974998L411.08438 30.881248Q411.08438 29.834373 411.27188 29.318748Q411.52188 28.631248 412.1625 28.209373Q412.80313 27.787498 413.95938 27.787498Q414.70938 27.787498 415.6 27.959373L415.33438 29.490623Q414.7875 29.396873 414.30313 29.396873Q413.50626 29.396873 413.1625 29.740623Q412.83438 30.084373 412.83438 31.021873L412.83438 31.974998L414.86563 31.974998L414.86563 33.35L412.83438 33.35L412.83438 42.35L411.08438 42.35ZM417.08438 42.35L417.08438 33.35L415.5375 33.35L415.5375 31.974998L417.08438 31.974998L417.08438 30.881248Q417.08438 29.834373 417.27188 29.318748Q417.52188 28.631248 418.1625 28.209373Q418.80313 27.787498 419.95938 27.787498Q420.70938 27.787498 421.6 27.959373L421.33438 29.490623Q420.7875 29.396873 420.30313 29.396873Q419.50626 29.396873 419.1625 29.740623Q418.83438 30.084373 418.83438 31.021873L418.83438 31.974998L420.86563 31.974998L420.86563 33.35L418.83438 33.35L418.83438 42.35L417.08438 42.35ZM429.77188 39.00625L431.58438 39.240623Q431.1625 40.81875 429.99063 41.709373Q428.83438 42.584373 427.02188 42.584373Q424.75626 42.584373 423.4125 41.19375Q422.08438 39.7875 422.08438 37.25625Q422.08438 34.63125 423.42813 33.19375Q424.77188 31.740623 426.92813 31.740623Q429.00626 31.740623 430.31876 33.1625Q431.64688 34.56875 431.64688 37.146873Q431.64688 37.303123 431.63126 37.615623L423.89688 37.615623Q423.99063 39.31875 424.86563 40.225Q425.74063 41.13125 427.0375 41.13125Q428.00626 41.13125 428.67813 40.63125Q429.36563 40.115623 429.77188 39.00625ZM423.99063 36.1625L429.7875 36.1625Q429.67813 34.865623 429.13126 34.209373Q428.2875 33.19375 426.94376 33.19375Q425.74063 33.19375 424.9125 34.00625Q424.08438 34.803123 423.99063 36.1625ZM433.64688 42.35L433.64688 31.974998L435.225 31.974998L435.225 33.553123Q435.83438 32.44375 436.35 32.1Q436.86563 31.740623 437.475 31.740623Q438.36563 31.740623 439.2875 32.303123L438.67813 33.94375Q438.0375 33.553123 437.39688 33.553123Q436.81876 33.553123 436.35 33.9125Q435.89688 34.25625 435.69376 34.865623Q435.4125 35.803123 435.4125 36.9125L435.4125 42.35L433.64688 42.35ZM447.77188 39.00625L449.58438 39.240623Q449.1625 40.81875 447.99063 41.709373Q446.83438 42.584373 445.02188 42.584373Q442.75626 42.584373 441.4125 41.19375Q440.08438 39.7875 440.08438 37.25625Q440.08438 34.63125 441.42813 33.19375Q442.77188 31.740623 444.92813 31.740623Q447.00626 31.740623 448.31876 33.1625Q449.64688 34.56875 449.64688 37.146873Q449.64688 37.303123 449.63126 37.615623L441.89688 37.615623Q441.99063 39.31875 442.86563 40.225Q443.74063 41.13125 445.0375 41.13125Q446.00626 41.13125 446.67813 40.63125Q447.36563 40.115623 447.77188 39.00625ZM441.99063 36.1625L447.7875 36.1625Q447.67813 34.865623 447.13126 34.209373Q446.2875 33.19375 444.94376 33.19375Q443.74063 33.19375 442.9125 34.00625Q442.08438 34.803123 441.99063 36.1625ZM451.6625 42.35L451.6625 31.974998L453.25626 31.974998L453.25626 33.459373Q454.39688 31.740623 456.55313 31.740623Q457.49063 31.740623 458.27188 32.084373Q459.06876 32.4125 459.45938 32.959373Q459.85 33.50625 460.00626 34.25625Q460.1 34.75625 460.1 35.975L460.1 42.35L458.33438 42.35L458.33438 36.0375Q458.33438 34.959373 458.13126 34.428123Q457.92813 33.896873 457.39688 33.584373Q456.88126 33.271873 456.17813 33.271873Q455.05313 33.271873 454.24063 33.990623Q453.42813 34.69375 453.42813 36.678123L453.42813 42.35L451.6625 42.35ZM465.50626 40.771873L465.75626 42.334373Q465.02188 42.490623 464.42813 42.490623Q463.475 42.490623 462.94376 42.19375Q462.42813 41.88125 462.20938 41.396873Q461.99063 40.896873 461.99063 39.31875L461.99063 33.35L460.70938 33.35L460.70938 31.974998L461.99063 31.974998L461.99063 29.412498L463.74063 28.349998L463.74063 31.974998L465.50626 31.974998L465.50626 33.35L463.74063 33.35L463.74063 39.4125Q463.74063 40.1625 463.83438 40.38125Q463.92813 40.584373 464.13126 40.725Q464.35 40.85 464.74063 40.85Q465.02188 40.85 465.50626 40.771873ZM472.31876 42.35L477.81876 28.037498L479.86563 28.037498L485.725 42.35L483.56876 42.35L481.89688 38.00625L475.9125 38.00625L474.33438 42.35L472.31876 42.35ZM476.45938 36.475L481.30313 36.475L479.81876 32.50625Q479.13126 30.693748 478.80313 29.537498Q478.52188 30.912498 478.02188 32.271873L476.45938 36.475ZM486.64688 42.35L486.64688 31.974998L488.225 31.974998L488.225 33.553123Q488.83438 32.44375 489.35 32.1Q489.86563 31.740623 490.475 31.740623Q491.36563 31.740623 492.2875 32.303123L491.67813 33.94375Q491.0375 33.553123 490.39688 33.553123Q489.81876 33.553123 489.35 33.9125Q488.89688 34.25625 488.69376 34.865623Q488.4125 35.803123 488.4125 36.9125L488.4125 42.35L486.64688 42.35ZM493.64688 42.35L493.64688 31.974998L495.225 31.974998L495.225 33.553123Q495.83438 32.44375 496.35 32.1Q496.86563 31.740623 497.475 31.740623Q498.36563 31.740623 499.2875 32.303123L498.67813 33.94375Q498.0375 33.553123 497.39688 33.553123Q496.81876 33.553123 496.35 33.9125Q495.89688 34.25625 495.69376 34.865623Q495.4125 35.803123 495.4125 36.9125L495.4125 42.35L493.64688 42.35ZM507.44376 41.06875Q506.45938 41.896873 505.55313 42.240623Q504.64688 42.584373 503.61563 42.584373Q501.9125 42.584373 500.99063 41.75625Q500.06876 40.9125 500.06876 39.615623Q500.06876 38.85 500.4125 38.225Q500.77188 37.6 501.33438 37.225Q501.89688 36.834373 502.6 36.63125Q503.11563 36.50625 504.1625 36.365623Q506.2875 36.115623 507.2875 35.771873Q507.30313 35.4125 507.30313 35.303123Q507.30313 34.240623 506.80313 33.7875Q506.13126 33.19375 504.80313 33.19375Q503.55313 33.19375 502.95938 33.63125Q502.38126 34.06875 502.1 35.178123L500.38126 34.94375Q500.61563 33.834373 501.14688 33.1625Q501.67813 32.475 502.69376 32.115623Q503.70938 31.740623 505.05313 31.740623Q506.38126 31.740623 507.20938 32.053123Q508.0375 32.365623 508.42813 32.85Q508.81876 33.31875 508.975 34.0375Q509.06876 34.490623 509.06876 35.6625L509.06876 38.00625Q509.06876 40.459373 509.17813 41.115623Q509.2875 41.75625 509.63126 42.35L507.7875 42.35Q507.52188 41.803123 507.44376 41.06875ZM507.2875 37.146873Q506.33438 37.5375 504.4125 37.803123Q503.33438 37.959373 502.88126 38.1625Q502.44376 38.35 502.19376 38.725Q501.94376 39.1 501.94376 39.56875Q501.94376 40.271873 502.475 40.740623Q503.00626 41.209373 504.0375 41.209373Q505.05313 41.209373 505.83438 40.771873Q506.63126 40.31875 507.00626 39.553123Q507.2875 38.94375 507.2875 37.7875L507.2875 37.146873ZM511.58438 46.35L511.39688 44.69375Q511.975 44.85 512.39685 44.85Q512.9906 44.85 513.33435 44.646873Q513.6937 44.459373 513.9125 44.115623Q514.08435 43.85 514.45935 42.803123Q514.5062 42.646873 514.6156 42.365623L510.67813 31.974998L512.5687 31.974998L514.725 37.990623Q515.14685 39.13125 515.475 40.38125Q515.7875 39.178123 516.1937 38.021873L518.4125 31.974998L520.1781 31.974998L516.225 42.521873Q515.6 44.240623 515.2406 44.88125Q514.77185 45.75625 514.1625 46.146873Q513.5687 46.553123 512.725 46.553123Q512.20935 46.553123 511.58438 46.35ZM527.2562 37.75625L529.0375 37.6Q529.1625 38.6625 529.6156 39.35Q530.08435 40.0375 531.0531 40.475Q532.0375 40.896873 533.27185 40.896873Q534.35 40.896873 535.1781 40.584373Q536.0062 40.25625 536.4125 39.69375Q536.8187 39.13125 536.8187 38.459373Q536.8187 37.7875 536.4281 37.2875Q536.0375 36.7875 535.14685 36.44375Q534.5687 36.209373 532.58435 35.740623Q530.6156 35.271873 529.83435 34.85Q528.8031 34.31875 528.3031 33.521873Q527.8031 32.725 527.8031 31.740623Q527.8031 30.646873 528.4125 29.709373Q529.02185 28.771873 530.20935 28.287498Q531.39685 27.787498 532.83435 27.787498Q534.4281 27.787498 535.64685 28.303123Q536.8656 28.818748 537.5062 29.818748Q538.1625 30.803123 538.20935 32.06875L536.39685 32.209373Q536.2562 30.849998 535.4125 30.162498Q534.5687 29.459373 532.9125 29.459373Q531.1937 29.459373 530.39685 30.099998Q529.6156 30.724998 529.6156 31.599998Q529.6156 32.38125 530.1781 32.88125Q530.725 33.38125 533.0375 33.896873Q535.35 34.4125 536.20935 34.803123Q537.45935 35.38125 538.0531 36.271873Q538.64685 37.146873 538.64685 38.303123Q538.64685 39.44375 537.9906 40.459373Q537.33435 41.475 536.1 42.0375Q534.8812 42.6 533.35 42.6Q531.4125 42.6 530.1 42.0375Q528.7875 41.459373 528.0375 40.31875Q527.2875 39.178123 527.2562 37.75625ZM540.6781 30.053123L540.6781 28.037498L542.4437 28.037498L542.4437 30.053123L540.6781 30.053123ZM540.6781 42.35L540.6781 31.974998L542.4437 31.974998L542.4437 42.35L540.6781 42.35ZM543.7406 42.35L543.7406 40.928123L550.35 33.35Q549.225 33.396873 548.3656 33.396873L544.1312 33.396873L544.1312 31.974998L552.6156 31.974998L552.6156 33.146873L546.9906 39.725L545.9125 40.928123Q547.08435 40.834373 548.1312 40.834373L552.9281 40.834373L552.9281 42.35L543.7406 42.35ZM560.77185 39.00625L562.58435 39.240623Q562.1625 40.81875 560.9906 41.709373Q559.83435 42.584373 558.02185 42.584373Q555.7562 42.584373 554.4125 41.19375Q553.08435 39.7875 553.08435 37.25625Q553.08435 34.63125 554.4281 33.19375Q555.77185 31.740623 557.9281 31.740623Q560.0062 31.740623 561.3187 33.1625Q562.64685 34.56875 562.64685 37.146873Q562.64685 37.303123 562.6312 37.615623L554.89685 37.615623Q554.9906 39.31875 555.8656 40.225Q556.7406 41.13125 558.0375 41.13125Q559.0062 41.13125 559.6781 40.63125Q560.3656 40.115623 560.77185 39.00625ZM554.9906 36.1625L560.7875 36.1625Q560.6781 34.865623 560.1312 34.209373Q559.2875 33.19375 557.9437 33.19375Q556.7406 33.19375 555.9125 34.00625Q555.08435 34.803123 554.9906 36.1625ZM563.95935 39.25625L565.70935 38.975Q565.85 40.021873 566.52185 40.584373Q567.1937 41.13125 568.39685 41.13125Q569.6 41.13125 570.1781 40.646873Q570.77185 40.146873 570.77185 39.490623Q570.77185 38.896873 570.2562 38.553123Q569.89685 38.31875 568.45935 37.959373Q566.52185 37.459373 565.77185 37.115623Q565.0375 36.75625 564.64685 36.13125Q564.2562 35.490623 564.2562 34.725Q564.2562 34.0375 564.5687 33.44375Q564.89685 32.85 565.4437 32.459373Q565.85 32.1625 566.5531 31.959373Q567.27185 31.740623 568.08435 31.740623Q569.3031 31.740623 570.225 32.1Q571.14685 32.44375 571.58435 33.053123Q572.02185 33.646873 572.1937 34.646873L570.475 34.88125Q570.35 34.084373 569.7875 33.646873Q569.225 33.19375 568.20935 33.19375Q566.9906 33.19375 566.475 33.6Q565.95935 33.990623 565.95935 34.521873Q565.95935 34.865623 566.1781 35.146873Q566.3812 35.428123 566.85 35.615623Q567.1156 35.709373 568.39685 36.053123Q570.27185 36.553123 571.0062 36.88125Q571.7406 37.19375 572.1625 37.803123Q572.58435 38.396873 572.58435 39.303123Q572.58435 40.178123 572.0687 40.959373Q571.5531 41.740623 570.58435 42.1625Q569.6156 42.584373 568.39685 42.584373Q566.3812 42.584373 565.3187 41.740623Q564.2562 40.896873 563.95935 39.25625Z" fill-rule="nonzero"/></svg>
\ No newline at end of file
diff --git a/src/main.cpp b/src/main.cpp
index 0abb168..3984f60 100644
--- a/src/main.cpp
+++ b/src/main.cpp
@@ -13,7 +13,7 @@
 #include <stream_compaction/thrust.h>
 #include "testing_helpers.hpp"
 
-const int SIZE = 1 << 12; // feel free to change the size of array = 256
+const int SIZE = 1 << 20; // feel free to change the size of array = 256
 const int NPOT = SIZE - 3; // Non-Power-Of-Two = 253
 int *a = new int[SIZE];
 int *b = new int[SIZE];

From dafa0e4ba4115c5dd09cc74b470994da4f18cf89 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 23:08:06 -0400
Subject: [PATCH 14/17] Update README.md

---
 README.md | 68 +++++++++++++++++++++++++++++++++----------------------
 1 file changed, 41 insertions(+), 27 deletions(-)

diff --git a/README.md b/README.md
index bcf05a1..bbe432c 100644
--- a/README.md
+++ b/README.md
@@ -42,10 +42,23 @@ Performance was pretty similar for most block sizes, but started to suffer for b
 | 67108864   | 100.736 | 38.8708  | 15.4414        | 2.14362  |
 | 268435456  | 410.365 | 169.486  | 60.6265        | 6.23341  |
 
+### Analysis
+The CPU implementation's run time appears to be linear with respect to the number of array elements. This makes sense because each element is processed one at a time inside a for loop.  
+
+Thrust is the fastest by far. This is probably because they are using shared memory, while all of my implementations only use global memory which is much slower to access, making each kernel thread slower. And maybe other optimizations as well.  
+
+The work-efficient scan is faster than the naive scan. This should be because I made optimizations (see next section) to reduce the number of threads at each iteration, whereas the naive scan still launches n threads each iteration.
+
+In all implementations, computation should not be a performance bottleneck, since each kernel runs in about O(1) time, we can't really do better than that. 
+
+Aside from the above trends, memory IO (cudaMemcpy) is a giant performance bottleneck. This is not shown in the performance graph since we start measuring runtime after the initial cudaMemcpy and stop measuring before the final cudaMemcpy. Still, cudaMemcpy runs in O(n) time, which effectively makes any GPU algorithm O(n), even though the actual implementation of scan runs in O(log n).  
+
+However, in practice, cudaMemcpy is still very fast, probably because the hardware bandwith for copying data from host to device is very large. For example, on an array of size 2^26, I ran my Radix sort algorithms and the CPU implementation took about 7 seconds (7131.1ms). Meanwhile the GPU implementation took about 1 second (826.585ms) including the cudaMemcpy, and half a second (434.71ms) without the cudaMemcpy. This means that while the cudaMemcpy is still a huge bottleneck on the GPU performance (taking up about half the runtime), it isn't too the point of being linear time, even at large numbers. In the future, I could try to measure the bandwidth of cudaMemcpy on my GPU.
+
 ### Extra credit
 
 #### Performance
-My work efficient scan halves the total number of threads launched each iteration, so the work efficient scan is faster at than naive and CPU implementation at array sizes of 2^18 and larger. This means less threads are idling and taking up space on the GPU multi-processors.
+My work efficient scan halves the total number of threads launched each iteration, this means less threads are idling and taking up space on the GPU multi-processors while other threads could be running. As a result, the work efficient scan is faster at than naive and CPU implementation at array sizes of 2^18 and larger.
 
 #### Radix sort
 
@@ -56,76 +69,77 @@ I implemented Radix sort on the CPU, wrote two test cases, and added Radix sort
 A few notes: you can pass in the number of bits you want to sort by, which should be `ilog2ceil(MAX_ARRAY_ELEMENT_VALUE)`. Also, I assumed for simplicity that each element is a positive integer (although still using int and not unsigned int types) so I can just use a bitmask to compact the arrays. Finally, to test, the array size should not be too close to 2^31 because of integer overflow issues...  
 
 #### Sample program output
-Settings: `blockSize` = 128, `SIZE` = 1 << 20
+Settings: `blockSize` = 128, `SIZE` = 1 << 26
 
 ```
 
 ****************
 ** SCAN TESTS **
 ****************
-    [   8  40  25  15  22   4  39  46  44  28   2  44  33 ...  30   0 ]
+    [  29  25   4  37   8  30  21  31   8  21  22  19   3 ...  49   0 ]
 ==== cpu scan, power-of-two ====
-   elapsed time: 1.5927ms    (std::chrono Measured)
-    [   0   8  48  73  88 110 114 153 199 243 271 273 317 ... 25674057 25674087 ]
+   elapsed time: 102.046ms    (std::chrono Measured)
+    [   0  29  54  58  95 103 133 154 185 193 214 236 255 ... 1643658734 1643658783 ]
 ==== cpu scan, non-power-of-two ====
-   elapsed time: 1.7177ms    (std::chrono Measured)
-    [   0   8  48  73  88 110 114 153 199 243 271 273 317 ... 25673959 25673964 ]
+   elapsed time: 99.084ms    (std::chrono Measured)
+    [   0  29  54  58  95 103 133 154 185 193 214 236 255 ... 1643658648 1643658670 ]
     passed
 ==== naive scan, power-of-two ====
-   elapsed time: 0.634144ms    (CUDA Measured)
+   elapsed time: 38.8846ms    (CUDA Measured)
     passed
 ==== naive scan, non-power-of-two ====
-   elapsed time: 0.602528ms    (CUDA Measured)
+   elapsed time: 37.4897ms    (CUDA Measured)
     passed
 ==== work-efficient scan, power-of-two ====
-   elapsed time: 0.457504ms    (CUDA Measured)
+   elapsed time: 15.4577ms    (CUDA Measured)
     passed
 ==== work-efficient scan, non-power-of-two ====
-   elapsed time: 0.479328ms    (CUDA Measured)
+   elapsed time: 15.4086ms    (CUDA Measured)
     passed
 ==== thrust scan, power-of-two ====
-   elapsed time: 0.340864ms    (CUDA Measured)
+   elapsed time: 2.09901ms    (CUDA Measured)
     passed
 ==== thrust scan, non-power-of-two ====
-   elapsed time: 0.279648ms    (CUDA Measured)
+   elapsed time: 2.36995ms    (CUDA Measured)
     passed
 
 *****************************
 ** STREAM COMPACTION TESTS **
 *****************************
-    [   2   0   2   2   3   1   2   1   1   2   2   0   3 ...   0   0 ]
+    [   0   3   2   3   3   1   3   2   2   3   1   3   0 ...   3   0 ]
 ==== cpu compact without scan, power-of-two ====
-   elapsed time: 2.2561ms    (std::chrono Measured)
-    [   2   2   2   3   1   2   1   1   2   2   3   1   2 ...   1   1 ]
+   elapsed time: 137.514ms    (std::chrono Measured)
+    [   3   2   3   3   1   3   2   2   3   1   3   1   3 ...   2   3 ]
     passed
 ==== cpu compact without scan, non-power-of-two ====
-   elapsed time: 2.1692ms    (std::chrono Measured)
-    [   2   2   2   3   1   2   1   1   2   2   3   1   2 ...   1   1 ]
+   elapsed time: 137.567ms    (std::chrono Measured)
+    [   3   2   3   3   1   3   2   2   3   1   3   1   3 ...   2   1 ]
     passed
 ==== cpu compact with scan ====
-   elapsed time: 5.498ms    (std::chrono Measured)
-    [   2   2   2   3   1   2   1   1   2   2   3   1   2 ...   1   1 ]
+   elapsed time: 348.893ms    (std::chrono Measured)
+    [   3   2   3   3   1   3   2   2   3   1   3   1   3 ...   2   3 ]
     passed
 ==== work-efficient compact, power-of-two ====
-   elapsed time: 0.541312ms    (CUDA Measured)
+   elapsed time: 19.1836ms    (CUDA Measured)
     passed
 ==== work-efficient compact, non-power-of-two ====
-   elapsed time: 0.490016ms    (CUDA Measured)
+   elapsed time: 19.1201ms    (CUDA Measured)
     passed
 
 *****************************
 ** RADIX SORT TESTS **
 *****************************
-    [ 23914 19768 8186 18274 8083 25457 19314 28109 22549 17382 25706 27432 8575 ... 4256   0 ]
+    [ 31399 13580 25635 22845 23360 14322 9628 3467 20074 16251 14385 30083 26014 ... 230   0 ]
 ==== cpu radix sort, power-of-two ====
+   elapsed time: 7131.1ms    (std::chrono Measured)
     [   0   0   0   0   0   0   0   0   0   0   0   0   0 ... 32767 32767 ]
 ==== radix sort, power of two ====
-   elapsed time: 0.490016ms    (CUDA Measured)
+   elapsed time: 826.585ms    (CUDA Measured)
     passed
 ==== cpu radix sort, non-power-of-two ====
-   elapsed time: 109.43ms    (std::chrono Measured)
+   elapsed time: 7102.31ms    (std::chrono Measured)
     [   0   0   0   0   0   0   0   0   0   0   0   0   0 ... 32767 32767 ]
 ==== radix sort, non-power of two ====
-   elapsed time: 0.490016ms    (CUDA Measured)
+   elapsed time: 788.974ms    (CUDA Measured)
     passed
-```
\ No newline at end of file
+```

From e7a4b08d7a9429813ffedeb1c21c885be02118b9 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 23:08:22 -0400
Subject: [PATCH 15/17] fix timing measurement

---
 src/main.cpp | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/src/main.cpp b/src/main.cpp
index 3984f60..6667cad 100644
--- a/src/main.cpp
+++ b/src/main.cpp
@@ -13,7 +13,7 @@
 #include <stream_compaction/thrust.h>
 #include "testing_helpers.hpp"
 
-const int SIZE = 1 << 20; // feel free to change the size of array = 256
+const int SIZE = 1 << 26; // feel free to change the size of array = 256
 const int NPOT = SIZE - 3; // Non-Power-Of-Two = 253
 int *a = new int[SIZE];
 int *b = new int[SIZE];
@@ -161,13 +161,13 @@ int main(int argc, char* argv[]) {
     zeroArray(SIZE, b);
     printDesc("cpu radix sort, power-of-two");
     StreamCompaction::CPU::radixSort(SIZE, RADIX_NUM_BITS, b, a);
-    //printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
+    printElapsedTime(StreamCompaction::CPU::timer().getCpuElapsedTimeForPreviousOperation(), "(std::chrono Measured)");
     printArray(SIZE, b, true);
 
     zeroArray(SIZE, c);
     printDesc("radix sort, power of two");
     StreamCompaction::Naive::radixSort(SIZE, RADIX_NUM_BITS, c, a);
-    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
+    printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
     printCmpResult(SIZE, b, c); // TODO: add cpu impl and write to b for comparison
 
     zeroArray(SIZE, b);
@@ -179,7 +179,7 @@ int main(int argc, char* argv[]) {
     zeroArray(SIZE, c);
     printDesc("radix sort, non-power of two");
     StreamCompaction::Naive::radixSort(NPOT, RADIX_NUM_BITS, c, a);
-    printElapsedTime(StreamCompaction::Efficient::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
+    printElapsedTime(StreamCompaction::Naive::timer().getGpuElapsedTimeForPreviousOperation(), "(CUDA Measured)");
     printCmpResult(NPOT, b, c); // TODO: add cpu impl and write to b for comparison
 
     system("pause"); // stop Win32 console from closing on exit

From 3036ae29140727a750312b19e336d991695b8524 Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 23:11:24 -0400
Subject: [PATCH 16/17] Update README.md

---
 README.md | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/README.md b/README.md
index bbe432c..0c56d3e 100644
--- a/README.md
+++ b/README.md
@@ -32,7 +32,7 @@ Performance was pretty similar for most block sizes, but started to suffer for b
 
 ![](img/Performance%20of%20Scan%20Implementations%20on%20Different%20Array%20Sizes.svg)
 
-| Array size | CPU (m/s)    | Naive (m/s)  | Work efficient (m/s) | Thrust (m/s)  |
+| Array size | CPU (ms)    | Naive (ms)  | Work efficient (ms) | Thrust (ms)  |
 |------------|---------|----------|----------------|----------|
 | 65536      | 0.1023  | 0.158464 | 0.266592       | 0.045472 |
 | 262144     | 0.399   | 0.2616   | 0.33888        | 0.194144 |

From 502d1b30839e0cd344df8bb9af2881bb148c8eae Mon Sep 17 00:00:00 2001
From: Constance Wang <constancezwang@gmail.com>
Date: Sun, 18 Sep 2022 23:13:14 -0400
Subject: [PATCH 17/17] Update README.md

---
 README.md | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/README.md b/README.md
index 0c56d3e..c0f022c 100644
--- a/README.md
+++ b/README.md
@@ -65,6 +65,8 @@ My work efficient scan halves the total number of threads launched each iteratio
 I implemented Radix sort on the CPU, wrote two test cases, and added Radix sort on the GPU, which calls my work efficient scan. The functions to look at are:
 - `radixSort` in `naive.cu`
 - `radixSort` in `cpu.cu`
+Example of usage:  
+`StreamCompaction::CPU::radixSort(NPOT, RADIX_NUM_BITS, b, a);`
 
 A few notes: you can pass in the number of bits you want to sort by, which should be `ilog2ceil(MAX_ARRAY_ELEMENT_VALUE)`. Also, I assumed for simplicity that each element is a positive integer (although still using int and not unsigned int types) so I can just use a bitmask to compact the arrays. Finally, to test, the array size should not be too close to 2^31 because of integer overflow issues...