Better Multithreading

.gitignore

@@ -3,3 +3,4 @@ cmake-build-debug
 *.txt
 merged
 *.out
+/build

main.cu

@@ -15,15 +15,32 @@
 #include <__clang_cuda_cmath.h>
 #endif
 
+// for windows plebs
+#ifdef __INTELLISENSE__
+
+#include <cuda.h>
+#include <cuda_runtime_api.h>
+#include <cuda_runtime.h>
+#define __CUDACC__ //fixes function defenition in ide
+//void __syncthreads();
+
+#include <device_launch_parameters.h>
+#include <device_functions.h>
+#include <device_atomic_functions.h>
+
+#endif
 
 
 #include <stdint.h>
 #include <memory.h>
 #include <stdio.h>
 #include <time.h>
 #include <ctype.h>
+
 #include <thread>
-#include <sstream>
+#include <vector>
+#include <mutex>
+#include <atomic>
 
 
 #define signed_seed_t int64_t
@@ -43,8 +60,8 @@
 #define RANDOM_SCALE 0x1p-48
 
 inline uint __host__ __device__  random_next(Random *random, int bits) {
-  *random = trunc((*random * RANDOM_MULTIPLIER + RANDOM_ADDEND) * RANDOM_SCALE);
-  return (uint)((ulong)(*random / RANDOM_SCALE) >> (48 - bits));
+    *random = trunc((*random * RANDOM_MULTIPLIER + RANDOM_ADDEND) * RANDOM_SCALE);
+    return (uint)((ulong)(*random / RANDOM_SCALE) >> (48 - bits));
 }
 
 #else
@@ -323,25 +340,7 @@ void calculate_search_backs() {
 #define OFFSET 0
 #endif
 
-void run(bool &running, int &offset, int &workSize, bool &finished, long long *&tempStorage, int &count, const std::string out_file_name) {
-    FILE* out_file = fopen(out_file_name.c_str(), "w");
-    while (running) {
-        if (workSize != 0 && !finished) {
-            for (int j = 0; j < workSize; ++j) {
-                if (generator::ChunkGenerator::populate(tempStorage[j + offset], X_TRANSLATE + 16)) {
-                    fprintf(out_file, "%lld\n", tempStorage[j]);
-                    count++;
-                }
-            }
-            finished = true;
-        }
-    }
-    fflush(out_file);
-    fclose(out_file);
-
-}
-
-int main(int argc, char *argv[]) {
+int main(int argc, char* argv[]) {
 #define int int32_t
     random_math::JavaRand::init();
     generator::ChunkGenerator::init();
@@ -351,34 +350,23 @@ int main(int argc, char *argv[]) {
 
     calculate_search_backs();
 
-    for(int i = 0; i < GPU_COUNT; i++) {
-        setup_gpu_node(&nodes[i],i);
-    }
+    FILE* out_file = fopen("chunk_seeds.txt", "w");
 
-    int threadCount = std::thread::hardware_concurrency();
-    threadCount = threadCount-4;
-    std::thread threads[threadCount];
-    long long *tempStorage = NULL;
-    bool running = true;
-    bool finishedArr[threadCount];
-    int offsetArr[threadCount];
-    int workSizeArr[threadCount];
-
-    int counts[threadCount];
-    printf("Using %d threads for cpu work\n",threadCount);
-    for (int i = 0; i < threadCount; ++i) {
-        counts[i] = 0;
-        finishedArr[i] = true;
-        offsetArr[i] = 0;
-        workSizeArr[i] = 0;
-        std::ostringstream oss;
-        oss << "chunk_seeds" << i << ".txt";
-        threads[i] = std::thread(run, std::ref(running), std::ref(offsetArr[i]), std::ref(workSizeArr[i]), std::ref(finishedArr[i]), std::ref(tempStorage), std::ref(counts[i]), oss.str());
+    for (int i = 0; i < GPU_COUNT; i++) {
+        setup_gpu_node(&nodes[i], i);
     }
-    ulong count = 0;
+
+    std::vector<std::thread> threads(std::thread::hardware_concurrency() - 4);
+    std::mutex fileMutex;
+
+    std::atomic<ulong> count(0);
     clock_t lastIteration = clock();
     clock_t startTime = clock();
+    long long* tempStorage = NULL;
     ulong arraySize = 0;
+
+    printf("Using %d threads for cpu work\n", (int)threads.size());
+
     for (ulong offset = OFFSET; offset < TOTAL_WORK_SIZE;) {
 
         for(int gpu_index = 0; gpu_index < GPU_COUNT; gpu_index++) {
@@ -401,62 +389,67 @@ int main(int argc, char *argv[]) {
             doWork<<<WORK_UNIT_SIZE / BLOCK_SIZE, BLOCK_SIZE>>>(nodes[gpu_index].num_tree_starts, nodes[gpu_index].tree_starts, nodes[gpu_index].num_seeds, nodes[gpu_index].seeds, search_back_count);
         }
 
+        static auto threadFunc = [&](size_t start, size_t end) {
+            for (int j = start; j < end; ++j) {
+                if (generator::ChunkGenerator::populate(tempStorage[j], X_TRANSLATE + 16)) {
+                    std::lock_guard<std::mutex> lock(fileMutex);
+                    fprintf(out_file, "%lld\n", tempStorage[j]);
+                    count++;
+                }
+            }
+        };
+
+
+        int chunkSize = arraySize / threads.size();
+        for(size_t i = 0; i < threads.size(); i++)
+            threads[i] = std::thread(threadFunc, i * chunkSize, (i == (threads.size() - 1)) ? arraySize : ((i + 1) * chunkSize));
+
+        for(std::thread& x : threads)
+            x.join();
+
+        fflush(out_file);
         free(tempStorage);
-        tempStorage = (long long*)malloc(sizeof(long long));
 
+        tempStorage = (long long*)malloc(sizeof(long long));
         arraySize = 0;
         for(int gpu_index = 0; gpu_index < GPU_COUNT; gpu_index++) {
             CHECK_GPU_ERR(cudaSetDevice(gpu_index));
             CHECK_GPU_ERR(cudaDeviceSynchronize());
             tempStorage = (long long*) realloc(tempStorage, (*nodes[gpu_index].num_seeds + arraySize) * sizeof(long long));
             for (int i = 0, e = *nodes[gpu_index].num_seeds; i < e; i++) {
                 tempStorage[arraySize + i] = nodes[gpu_index].seeds[i];
+                //fprintf(out_file, "%lld\n", nodes[gpu_index].seeds[i]);
             }
+            //fflush(out_file);
             arraySize += *nodes[gpu_index].num_seeds;
         }
 
-        int workSize = arraySize / threadCount;
-
-        if (workSize != 0) {
-            for (int k = 0; arraySize != 0 && k < threadCount; ++k) {
-                offsetArr[k] = k * workSize;
-                workSizeArr[k] = k != threadCount - 1 ? workSize : arraySize % workSize;
-                finishedArr[k] = workSizeArr[k] == 0;
-            }
+        double iterationTime = (double)(clock() - lastIteration) / CLOCKS_PER_SEC;
+        double timeElapsed = (double)(clock() - startTime) / CLOCKS_PER_SEC;
+        lastIteration = clock();
+        ulong numSearched = offset + WORK_UNIT_SIZE * GPU_COUNT - OFFSET;
+        double speed = numSearched / timeElapsed / 1000000;
+        double progress = (double)numSearched / (double)TOTAL_WORK_SIZE * 100.0;
+        double estimatedTime = (double)(TOTAL_WORK_SIZE - numSearched) / speed / 1000000;
+        char suffix = 's';
+        if (estimatedTime >= 3600) {
+            suffix = 'h';
+            estimatedTime /= 3600.0;
         }
-
-        if (workSize != 0) {
-
-            for (int j = 0; j < threadCount; ++j) {
-                count += counts[j];
-            }
-            double iterationTime = (double)(clock() - lastIteration) / CLOCKS_PER_SEC;
-            double timeElapsed = (double)(clock() - startTime) / CLOCKS_PER_SEC;
-            lastIteration = clock();
-            ulong numSearched = offset + WORK_UNIT_SIZE * GPU_COUNT - OFFSET;
-            double speed = numSearched / timeElapsed / 1000000;
-            double progress = (double)numSearched / (double)TOTAL_WORK_SIZE * 100.0;
-            double estimatedTime = (double)(TOTAL_WORK_SIZE - numSearched) / speed / 1000000;
-            char suffix = 's';
-            if (estimatedTime >= 3600) {
-                suffix = 'h';
-                estimatedTime /= 3600.0;
-            } else if (estimatedTime >= 60) {
-                suffix = 'm';
-                estimatedTime /= 60.0;
-            }
-            if (progress >= 100.0) {
-                estimatedTime = 0.0;
-                suffix = 's';
-            }
-            printf("Searched: %ld seeds. Found: %ld matches. Uptime: %.1fs. Speed: %.2fm seeds/s. Completion: %.3f%%. ETA: %.1f%c.\n", numSearched, count, timeElapsed, speed, progress, estimatedTime, suffix);
-            count = 0;
+        else if (estimatedTime >= 60) {
+            suffix = 'm';
+            estimatedTime /= 60.0;
         }
+        if (progress >= 100.0) {
+            estimatedTime = 0.0;
+            suffix = 's';
+        }
+        printf("Searched: %ld seeds. Found: %ld matches. Uptime: %.1fs. Speed: %.2fm seeds/s. Completion: %.3f%%. ETA: %.1f%c.\n", numSearched, count.load(), timeElapsed, speed, progress, estimatedTime, suffix);
+
     }
 
-    // Last batch to do
-    FILE* out_file = fopen("chunk_seeds_last.txt", "w");
 
+    // Last batch to do
     for (int j = 0; j < arraySize; ++j) {
         if (generator::ChunkGenerator::populate(tempStorage[j], X_TRANSLATE + 16)) {
             fprintf(out_file, "%lld\n", tempStorage[j]);
@@ -468,4 +461,4 @@ int main(int argc, char *argv[]) {
     free(tempStorage);
     fclose(out_file);
 
-}
+}