Second commit

CMakeLists.txt

@@ -9,11 +9,20 @@ find_package(SDL2 REQUIRED)
 
 add_subdirectory(glm)
 
-set(CMAKE_CUDA_FLAGS ${CMAKE_CUDA_FLAGS} "-G")
+set(CMAKE_CUDA_FLAGS ${CMAKE_CUDA_FLAGS} "")
 
-add_executable(cudagravity main.cu particle.cu particle.cuh definitions.cuh cudaMemory.cuh sdlHelper.cu sdlHelper.cuh)
+# Simulator
 
-target_link_libraries(cudagravity PUBLIC glm::glm SDL2::SDL2)
+add_executable(cudagravity_sim sim/main.cu common/particle.hpp common/definitions.hpp sim/cudaMemory.cuh sim/exportHelper.cu sim/exportHelper.cuh)
 
-set_target_properties(cudagravity PROPERTIES
+target_include_directories(cudagravity_sim PUBLIC common)
+
+target_link_libraries(cudagravity_sim PUBLIC glm::glm SDL2::SDL2)
+
+set_target_properties(cudagravity_sim PROPERTIES
         CUDA_SEPARABLE_COMPILATION ON)
+
+# Viewer
+
+add_executable(cudagravity_viewer viewer/main.cpp common/particle.hpp common/definitions.hpp)
+target_include_directories(cudagravity_viewer PUBLIC common)

common/definitions.hpp

@@ -0,0 +1,3 @@
+#include "glm/vec2.hpp"
+
+typedef glm::vec<2, float, glm::defaultp> vec2;

common/particle.hpp

@@ -0,0 +1,14 @@
+#ifndef VERLET_PARTICLE_HPP
+#define VERLET_PARTICLE_HPP
+
+#include "definitions.hpp"
+
+class Particle {
+public:
+    float mass;
+    vec2 position;
+    vec2 velocity;
+};
+
+
+#endif //VERLET_PARTICLE_HPP

definitions.cuh

@@ -1,4 +0,0 @@
-#include <glm/vec2.hpp>
-
-typedef glm::vec<2, float, glm::defaultp> vec2;
-#define CUDA_CALL(CALL) if (CALL != cudaSuccess) { printf("CUDA call failed at %s:%i\n", __FILE__,__LINE__); exit(-1); }

particle.cu

@@ -1,5 +0,0 @@
-//
-// Created by kai on 02/02/23.
-//
-
-#include "particle.cuh"

particle.cuh

@@ -1,16 +0,0 @@
-
-
-#ifndef VERLET_PARTICLE_CUH
-#define VERLET_PARTICLE_CUH
-
-#include "definitions.cuh"
-
-class Particle {
-public:
-    float mass;
-    vec2 position;
-    vec2 velocity;
-};
-
-
-#endif //VERLET_PARTICLE_CUH

sdlHelper.cu

@@ -1,88 +0,0 @@
-//
-// Created by kai on 05/02/23.
-//
-
-#include "sdlHelper.cuh"
-
-sdlHelper::sdlHelper() {
-    if (getenv("NO_DISPLAY") != nullptr) {
-        this->displaying = false;
-        return;
-    }
-
-    if (instance == nullptr) {
-        instance = this;
-
-        this->renderThread = std::thread(&sdlHelper::renderThreadFunc, this);
-    }
-}
-
-sdlHelper::~sdlHelper() {
-    if (instance == this) {
-        running = false;
-
-        this->renderThread.join();
-        SDL_Quit();
-    }
-}
-
-bool sdlHelper::drawParticles(Particle* particles, size_t particleCount) {
-    if (!this->displaying) return this->running;
-    this->particleRects.clear();
-    this->particleRects.reserve(particleCount);
-
-    float maxDistance = 0;
-
-    for (int i = 0; i < particleCount; i++) {
-        maxDistance = std::max(maxDistance, particles[i].position.x);
-        maxDistance = std::max(maxDistance, particles[i].position.y);
-    }
-
-    for (int i = 0; i < particleCount; i++) {
-        SDL_Rect rect;
-
-        vec2 pos = particles[i].position;
-        pos /= (maxDistance * 1.1);
-        pos *= 512;
-        pos += vec2(512, 512);
-
-        rect.x = pos.x - 1;
-        rect.y = pos.y - 1;
-        rect.w = 2;
-        rect.h = 2;
-        this->particleRects.push_back(rect);
-    }
-    return this->running;
-}
-
-sdlHelper *sdlHelper::getInstance() {
-    return this->instance;
-}
-
-void sdlHelper::renderThreadFunc() {
-    this->window = SDL_CreateWindow("cudaGravity", SDL_WINDOWPOS_CENTERED, SDL_WINDOWPOS_CENTERED, 1024, 1024, SDL_WINDOW_VULKAN);
-    this->renderer = SDL_CreateRenderer(this->window, -1, SDL_RENDERER_ACCELERATED | SDL_RENDERER_PRESENTVSYNC);
-
-    while (running) {
-        SDL_Event e;
-
-        while (SDL_PollEvent(&e)) {
-            if (e.type == SDL_QUIT) {
-                running = false;
-            }
-        }
-
-        SDL_SetRenderDrawColor(this->renderer, 0, 0, 0, 255);
-        SDL_RenderClear(this->renderer);
-        SDL_SetRenderDrawColor(this->renderer, 255, 255, 255, 255);
-
-        for (auto & particleRect : particleRects) {
-            SDL_RenderFillRect(this->renderer, &particleRect);
-        }
-
-        SDL_RenderPresent(this->renderer);
-    }
-
-    SDL_DestroyRenderer(this->renderer);
-    SDL_DestroyWindow(this->window);
-}

sdlHelper.cuh

@@ -1,35 +0,0 @@
-//
-// Created by kai on 05/02/23.
-//
-
-#ifndef VERLET_SDLHELPER_CUH
-#define VERLET_SDLHELPER_CUH
-
-#include <SDL2/SDL.h>
-#include <vector>
-#include <thread>
-#include "particle.cuh"
-
-class sdlHelper {
-public:
-    sdlHelper();
-    ~sdlHelper();
-
-    bool drawParticles(Particle* particles, size_t particleCount);
-    sdlHelper* getInstance();
-private:
-    void renderThreadFunc();
-
-    sdlHelper* instance = nullptr;
-    SDL_Window* window = nullptr;
-    SDL_Renderer* renderer = nullptr;
-    bool running = true;
-    bool displaying = true;
-
-    std::thread renderThread;
-
-    std::vector<SDL_Rect> particleRects;
-};
-
-
-#endif //VERLET_SDLHELPER_CUH

sim/cudaMemory.cuh

@@ -7,9 +7,10 @@
 #ifndef VERLET_CUDAMEMORY_CUH
 #define VERLET_CUDAMEMORY_CUH
 
-#include "definitions.cuh"
+#include "../common/definitions.hpp"
 #include <cstdio>
 
+#define CUDA_CALL(CALL) if (CALL != cudaSuccess) { printf("CUDA call failed at %s:%i\n", __FILE__,__LINE__); exit(-1); }
 
 template<typename T>
 class CudaMemory {
@@ -46,13 +47,25 @@ public:
     }
 
     void sync() {
-        CUDA_CALL(cudaMemcpy(this->hostPtr, this->devPtr, sizeof(T) * this->N, cudaMemcpyDeviceToHost));
+        CUDA_CALL(cudaMemcpyAsync(this->hostPtr, this->devPtr, sizeof(T) * this->N, cudaMemcpyDeviceToHost));
+    }
+
+    void sync_to_async(T** ptr) {
+        CUDA_CALL(cudaMemcpyAsync(*ptr, this->devPtr, sizeof(T) * this->N, cudaMemcpyDeviceToHost));
+    }
+
+    void sync_to(T** ptr) {
+        CUDA_CALL(cudaMemcpy(*ptr, this->devPtr, sizeof(T) * this->N, cudaMemcpyDeviceToHost));
     }
 
     T& operator[](size_t idx) {
         return this->hostPtr[idx];
     }
 
+    size_t size() {
+        return this->N * sizeof(T);
+    }
+
 private:
     T* hostPtr;
     T* devPtr;

sim/exportHelper.cu

@@ -0,0 +1,94 @@
+//
+// Created by kai on 05/02/23.
+//
+
+#include "exportHelper.cuh"
+
+#include <sys/stat.h>
+
+exportHelper::exportHelper() {
+    mkdir("output", S_IRWXU);
+
+    char* epochs_str = getenv("EPOCHS");
+
+    if (epochs_str != nullptr) {
+        if (sscanf(epochs_str, "%zu", &this->n) == 0) {
+            printf("Failed to parse EPOCHS environment variable - required when NO_DISPLAY is set\n");
+            should_stop = false;
+        }
+    } else {
+        printf("Failed to parse EPOCHS environment variable - required when NO_DISPLAY is set\n");
+        should_stop = false;
+    }
+
+    if (instance == nullptr) {
+        instance = this;
+    }
+}
+
+void exportHelper::stop() {
+    if (instance == this) {
+        should_stop = false;
+
+        for (auto& thread : exportThreads) {
+            if (thread.running) {
+                thread.thread.join();
+            }
+        }
+
+    }
+}
+
+exportHelper::~exportHelper() {
+    stop();
+}
+
+exportHelper *exportHelper::getInstance() {
+    return this->instance;
+}
+
+void exportHelper::imageWriter(size_t idx, thread* thread) {
+    char filename[256];
+    sprintf(filename, "output/%zu.dat", idx);
+
+    FILE* f = fopen64(filename, "wb");
+
+    auto* particles_copy = (Particle*)malloc(this->particles->size());
+    this->particles->sync_to(&particles_copy);
+
+    for (size_t i = 0; i < particle_count; i++) {
+        fprintf(f, "%f %f %f %f %f\n", particles_copy[i].position.x, particles_copy[i].position.y, particles_copy[i].velocity.x, particles_copy[i].velocity.y, particles_copy[i].mass);
+    }
+
+    free(particles_copy);
+    fflush(f);
+    fclose(f);
+
+    thread->running = false;
+}
+
+void exportHelper::epoch() {
+    this->count++;
+
+    if (this->count % 100 == 0) {
+        printf("Epoch: %zu\n", this->count);
+    }
+
+    if (this->n != 0 && this->count >= this->n) {
+        this->should_stop = false;
+    }
+
+    this->exportThreads.emplace_back();
+
+    this->exportThreads.back().thread = std::thread(&exportHelper::imageWriter, this, this->count, &this->exportThreads.back());
+
+    for (auto it = this->exportThreads.begin(); it != this->exportThreads.end();) {
+        if (!it->running) {
+            it->thread.join();
+            it = this->exportThreads.erase(it);
+        } else {
+            it++;
+        }
+    }
+
+}

sim/exportHelper.cuh

@@ -0,0 +1,53 @@
+//
+// Created by kai on 05/02/23.
+//
+
+#ifndef VERLET_EXPORTHELPER_CUH
+#define VERLET_EXPORTHELPER_CUH
+
+#include <vector>
+#include <thread>
+#include <list>
+#include <atomic>
+#include <particle.hpp>
+#include "cudaMemory.cuh"
+
+struct thread {
+    std::atomic<bool> running = true;
+    std::thread thread = std::thread();
+};
+
+class exportHelper {
+public:
+    bool should_stop = true;
+
+    exportHelper();
+    ~exportHelper();
+
+    exportHelper* getInstance();
+
+    void setParticles(CudaMemory<Particle>* _particles, size_t N) {
+        this->particles = _particles;
+        this->particle_count = N;
+    }
+
+    void epoch();
+
+    void stop();
+private:
+    exportHelper* instance = nullptr;
+    bool displaying = true;
+
+    size_t n = 0;
+    size_t count = 0;
+
+    CudaMemory<Particle>* particles;
+    size_t particle_count;
+
+    std::list<thread> exportThreads;
+
+    void imageWriter(size_t idx, thread* thread);
+};
+
+
+#endif //VERLET_EXPORTHELPER_CUH

sim/main.cu

@@ -1,13 +1,13 @@
 #include <iostream>
-#include "definitions.cuh"
-#include "particle.cuh"
+#include "../common/definitions.hpp"
+#include "../common/particle.hpp"
 #include "cudaMemory.cuh"
-#include "sdlHelper.cuh"
-#include <glm/glm.hpp>
-#include <random>
+#include "exportHelper.cuh"
+#include "glm/glm.hpp"
+#include <unistd.h>
 
 __constant__ const float G_CONSTANT = 1;
-__constant__ float dt = 0.1;
+__constant__ float dt = 0.0001;
 __constant__ const float softening = 0.0000;
 __constant__ const dim3 threadsPerBlock_gpu = dim3(1024);
 const dim3 threadsPerBlock_cpu = dim3(1024);
@@ -23,19 +23,14 @@ __global__ void getAccels(Particle* particles, vec2* accels, size_t N, size_t i)
             return;
         }
 
-        Particle particle = particles[i];
-        Particle other_particle = particles[j];
+        Particle p2 = particles[i];
+        Particle p1 = particles[j];
 
-        vec2 d = {
-                other_particle.position.x - particle.position.x,
-                other_particle.position.y - particle.position.y
-        };
+        float dist = glm::distance(p1.position, p2.position);
 
-        float dist = glm::distance(particle.position, other_particle.position);
+        vec2 f = -G_CONSTANT * ((p2.mass * p1.mass) / (dist * dist)) * glm::normalize(p2.position - p1.position);
 
-        vec2 f = -G_CONSTANT * ((particle.mass * other_particle.mass) / (dist * dist)) * glm::normalize(other_particle.position - particle.position);
-
-        vec2 accel = f / particle.mass;
+        vec2 accel = f / p2.mass;
 
         atomicAdd(&(accels[i].x), accel.x);
         atomicAdd(&(accels[i].y), accel.y);
@@ -63,29 +58,26 @@ __global__ void run_step(Particle* particles, vec2* accels, size_t N) {
     }
 }
 
-int main() {
-    sdlHelper sdlHelper;
+float get_random() {
+    float r = arc4random();
 
-    size_t N = 20000;
+    return (r / ((float)UINT32_MAX / 2.0f)) - 1;
+}
+
+int main() {
+    exportHelper sdlHelper;
+
+    size_t N = 200;
 
     CudaMemory<Particle> particles = CudaMemory<Particle>(N);
     CudaMemory<vec2> accelerations = CudaMemory<vec2>(N);
 
-    std::default_random_engine generator(std::random_device{}());
-    std::uniform_real_distribution<float> mass_distribution(1.0, 100000.0);
-    std::uniform_real_distribution<float> pos_distribution(-0.1, 0.1);
-    std::uniform_real_distribution<float> vel_distribution(-2.0, 2.0);
-
-    for (size_t i = 0; i < N; i++) {
-        accelerations[i] = {0, 0};
-        particles[i] = {
-                mass_distribution(generator),
-                {pos_distribution(generator), pos_distribution(generator)},
-                {vel_distribution(generator), vel_distribution(generator)}
-        };
+    for (int i = 0; i < N; i++) {
+        accelerations[i] = vec2(0, 0);
+        particles[i] = Particle{1, vec2(get_random(), get_random()), vec2(get_random(), get_random())};
     }
 
-    cudaDeviceGetAttribute((int*)(&threadsPerBlock_cpu.x), cudaDevAttrMaxThreadsPerBlock, 0);
+    //cudaDeviceGetAttribute((int*)(&threadsPerBlock_cpu.x), cudaDevAttrMaxThreadsPerBlock, 0);
 
     dim3 n_blocksPerGrid = dim3(
             ((N) + threadsPerBlock_cpu.x - 1) / threadsPerBlock_cpu.x
@@ -100,18 +92,24 @@ int main() {
     particles.send();
     accelerations.send();
 
-    for (int i = 0; i < 2; i++) {
+    sdlHelper.setParticles(&particles, N);
+
+    while(true) {
         run_step<<<n_blocksPerGrid, threadsPerBlock_cpu>>>(particles.getDevicePointer(), accelerations.getDevicePointer(), N);
 
         cudaDeviceSynchronize();
-        particles.sync();
+
         accelerations.send();
 
-        if (!sdlHelper.drawParticles(particles.getPointer(), N)) {
+        sdlHelper.epoch();
+
+        if (!sdlHelper.should_stop) {
             break;
         }
     }
 
+    sdlHelper.stop();
+
     cudaDeviceSynchronize();
 
     return 0;

viewer/main.cpp

@@ -0,0 +1,7 @@
+//
+// Created by kai on 07/02/23.
+//
+
+int main() {
+    return 0;
+}