diff --git a/CMakeLists.txt b/CMakeLists.txt index c8decc3..2dbb3e0 100644 --- a/CMakeLists.txt +++ b/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) \ No newline at end of file diff --git a/common/definitions.hpp b/common/definitions.hpp new file mode 100644 index 0000000..f0ffcf2 --- /dev/null +++ b/common/definitions.hpp @@ -0,0 +1,3 @@ +#include "glm/vec2.hpp" + +typedef glm::vec<2, float, glm::defaultp> vec2; diff --git a/common/particle.hpp b/common/particle.hpp new file mode 100644 index 0000000..a23a36c --- /dev/null +++ b/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 diff --git a/definitions.cuh b/definitions.cuh deleted file mode 100644 index 02b01b2..0000000 --- a/definitions.cuh +++ /dev/null @@ -1,4 +0,0 @@ -#include - -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); } diff --git a/objects.def b/objects.def deleted file mode 100644 index e69de29..0000000 diff --git a/particle.cu b/particle.cu deleted file mode 100644 index a863e80..0000000 --- a/particle.cu +++ /dev/null @@ -1,5 +0,0 @@ -// -// Created by kai on 02/02/23. -// - -#include "particle.cuh" diff --git a/particle.cuh b/particle.cuh deleted file mode 100644 index cbcb151..0000000 --- a/particle.cuh +++ /dev/null @@ -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 diff --git a/sample_data/random.7z b/sample_data/random.7z new file mode 100644 index 0000000..4aef09e Binary files /dev/null and b/sample_data/random.7z differ diff --git a/sdlHelper.cu b/sdlHelper.cu deleted file mode 100644 index 74f1ff5..0000000 --- a/sdlHelper.cu +++ /dev/null @@ -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); -} diff --git a/sdlHelper.cuh b/sdlHelper.cuh deleted file mode 100644 index e7b1f5f..0000000 --- a/sdlHelper.cuh +++ /dev/null @@ -1,35 +0,0 @@ -// -// Created by kai on 05/02/23. -// - -#ifndef VERLET_SDLHELPER_CUH -#define VERLET_SDLHELPER_CUH - -#include -#include -#include -#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 particleRects; -}; - - -#endif //VERLET_SDLHELPER_CUH diff --git a/cudaMemory.cuh b/sim/cudaMemory.cuh similarity index 66% rename from cudaMemory.cuh rename to sim/cudaMemory.cuh index a0052f3..3e870e1 100644 --- a/cudaMemory.cuh +++ b/sim/cudaMemory.cuh @@ -7,9 +7,10 @@ #ifndef VERLET_CUDAMEMORY_CUH #define VERLET_CUDAMEMORY_CUH -#include "definitions.cuh" +#include "../common/definitions.hpp" #include +#define CUDA_CALL(CALL) if (CALL != cudaSuccess) { printf("CUDA call failed at %s:%i\n", __FILE__,__LINE__); exit(-1); } template 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; diff --git a/sim/exportHelper.cu b/sim/exportHelper.cu new file mode 100644 index 0000000..e8aa0f7 --- /dev/null +++ b/sim/exportHelper.cu @@ -0,0 +1,94 @@ +// +// Created by kai on 05/02/23. +// + +#include "exportHelper.cuh" + +#include + +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++; + } + } + +} \ No newline at end of file diff --git a/sim/exportHelper.cuh b/sim/exportHelper.cuh new file mode 100644 index 0000000..9992d29 --- /dev/null +++ b/sim/exportHelper.cuh @@ -0,0 +1,53 @@ +// +// Created by kai on 05/02/23. +// + +#ifndef VERLET_EXPORTHELPER_CUH +#define VERLET_EXPORTHELPER_CUH + +#include +#include +#include +#include +#include +#include "cudaMemory.cuh" + +struct thread { + std::atomic running = true; + std::thread thread = std::thread(); +}; + +class exportHelper { +public: + bool should_stop = true; + + exportHelper(); + ~exportHelper(); + + exportHelper* getInstance(); + + void setParticles(CudaMemory* _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* particles; + size_t particle_count; + + std::list exportThreads; + + void imageWriter(size_t idx, thread* thread); +}; + + +#endif //VERLET_EXPORTHELPER_CUH diff --git a/main.cu b/sim/main.cu similarity index 58% rename from main.cu rename to sim/main.cu index eba49e2..e266ebb 100644 --- a/main.cu +++ b/sim/main.cu @@ -1,13 +1,13 @@ #include -#include "definitions.cuh" -#include "particle.cuh" +#include "../common/definitions.hpp" +#include "../common/particle.hpp" #include "cudaMemory.cuh" -#include "sdlHelper.cuh" -#include -#include +#include "exportHelper.cuh" +#include "glm/glm.hpp" +#include __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 particles = CudaMemory(N); CudaMemory accelerations = CudaMemory(N); - std::default_random_engine generator(std::random_device{}()); - std::uniform_real_distribution mass_distribution(1.0, 100000.0); - std::uniform_real_distribution pos_distribution(-0.1, 0.1); - std::uniform_real_distribution 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<<>>(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; diff --git a/viewer/main.cpp b/viewer/main.cpp new file mode 100644 index 0000000..9d8247a --- /dev/null +++ b/viewer/main.cpp @@ -0,0 +1,7 @@ +// +// Created by kai on 07/02/23. +// + +int main() { + return 0; +} \ No newline at end of file