Extract first public LiterateLB version

7 files changed, 423 insertions, 0 deletions

diff --git a/tangle/util/camera.h b/tangle/util/camera.h
new file mode 100644
index 0000000..0a793d3
--- /dev/null
+++ b/tangle/util/camera.h
@@ -0,0 +1,86 @@
+#include <cuda-samples/Common/helper_math.h>
+#include "SFML/Window/Event.hpp"
+
+class Camera {
+private:
+  float _distance;
+  float _phi;
+  float _psi;
+  float3 _target;
+  float3 _eye;
+  float3 _direction;
+  bool _dragging;
+  bool _moving;
+  float2 _lastMouse;
+
+public:
+  Camera(float3 target, float distance, float phi, float psi):
+    _distance(distance),
+    _phi(phi),
+    _psi(psi),
+    _target(target),
+    _dragging(false),
+    _moving(false) {
+    update();
+  }
+
+  void update() {
+    _eye = _target + make_float3(_distance*sin(_psi)*cos(_phi), _distance*sin(_psi)*sin(_phi), _distance*cos(_psi));
+    _direction = normalize(_target - _eye);
+  }
+
+  void handle(sf::Event& event) {
+    switch (event.type) {
+    case sf::Event::MouseWheelMoved:
+      _distance -= event.mouseWheel.delta * 10;
+      break;
+    case sf::Event::MouseButtonPressed:
+      if (event.mouseButton.button == sf::Mouse::Left) {
+        _dragging = true;
+        _lastMouse = make_float2(event.mouseButton.x, event.mouseButton.y);
+      } else if (event.mouseButton.button == sf::Mouse::Right) {
+        _moving = true;
+        _lastMouse = make_float2(event.mouseButton.x, event.mouseButton.y);
+      }
+      break;
+    case sf::Event::MouseButtonReleased:
+      if (event.mouseButton.button == sf::Mouse::Left) {
+        _dragging = false;
+      } else if (event.mouseButton.button == sf::Mouse::Right) {
+        _moving = false;
+      }
+      break;
+    case sf::Event::MouseMoved:
+      if (_dragging) {
+        float2 mouse = make_float2(event.mouseMove.x, event.mouseMove.y);
+        float2 delta = mouse - _lastMouse;
+        _lastMouse = mouse;
+        _phi += 0.4*delta.x * 2*M_PI/360;
+        if (delta.y > 0 && _psi <= M_PI-2*M_PI/60) {
+          _psi += 0.4*delta.y * M_PI/180;
+        } else if (delta.y < 0 && _psi >= 2*M_PI/60) {
+          _psi += 0.4*delta.y * M_PI/180;
+        }
+      }
+      if (_moving) {
+        float2 mouse = make_float2(event.mouseMove.x, event.mouseMove.y);
+        float2 delta = mouse - _lastMouse;
+        _lastMouse = mouse;
+        float3 forward = normalize(_target - _eye);
+    	   float3 right   = normalize(cross(make_float3(0.f, 0.f, -1.f), forward));
+        float3 up      = cross(right, forward);
+        _target += 0.4*right*delta.x - 0.4*up*delta.y;
+      }
+      break;
+    }
+    update();
+  }
+
+  float3 getDirection() const {
+    return _direction;
+  }
+
+  float3 getEyePosition() const {
+    return _eye;
+  }
+};
diff --git a/tangle/util/colormap.h b/tangle/util/colormap.h
new file mode 100644
index 0000000..c01cea8
--- /dev/null
+++ b/tangle/util/colormap.h
@@ -0,0 +1,38 @@
+#pragma once
+#include "assets.h"
+#include "texture.h"
+
+#include <imgui.h>
+#include <imgui-SFML.h>
+#include <SFML/Graphics.hpp>
+
+struct ColorPalette {
+  const assets::File* current;
+  sf::Texture texture;
+
+  ColorPalette(cudaSurfaceObject_t& palette) {
+    current = &assets::palette::files[5];
+    texture.loadFromMemory(current->data, current->size);
+    palette = bindTextureToCuda(texture);
+  }
+
+  void interact();
+};
+
+void ColorPalette::interact() {
+  if (ImGui::BeginCombo("Color palette", current->name.c_str())) {
+    for (unsigned i=0; i < assets::palette::file_count; ++i) {
+      bool is_selected = (current == &assets::palette::files[i]);
+      if (ImGui::Selectable(assets::palette::files[i].name.c_str(), is_selected)) {
+        current = &assets::palette::files[i];
+        texture.loadFromMemory(current->data, current->size);
+        break;
+      }
+      if (is_selected) {
+        ImGui::SetItemDefaultFocus();
+      }
+    }
+    ImGui::EndCombo();
+  }
+  ImGui::Image(texture, sf::Vector2f(400.,40.));
+}
diff --git a/tangle/util/noise.h b/tangle/util/noise.h
new file mode 100644
index 0000000..22ad548
--- /dev/null
+++ b/tangle/util/noise.h
@@ -0,0 +1,39 @@
+#pragma once
+#include "assets.h"
+#include "texture.h"
+
+#include <imgui.h>
+#include <imgui-SFML.h>
+#include <SFML/Graphics.hpp>
+
+struct NoiseSource {
+  const assets::File* current;
+  sf::Texture texture;
+
+  NoiseSource(cudaSurfaceObject_t& noise) {
+    current = &assets::noise::files[0];
+    texture.loadFromMemory(current->data, current->size);
+    noise = bindTextureToCuda(texture);
+  }
+
+  void interact();
+};
+
+void NoiseSource::interact() {
+  ImGui::Image(texture, sf::Vector2f(32,20));
+  ImGui::SameLine();
+  if (ImGui::BeginCombo("Noise", current->name.c_str())) {
+    for (unsigned i=0; i < assets::noise::file_count; ++i) {
+      bool is_selected = (current == &assets::noise::files[i]);
+      if (ImGui::Selectable(assets::noise::files[i].name.c_str(), is_selected)) {
+        current = &assets::noise::files[i];
+        texture.loadFromMemory(current->data, current->size);
+        break;
+      }
+      if (is_selected) {
+        ImGui::SetItemDefaultFocus();
+      }
+    }
+    ImGui::EndCombo();
+  }
+}
diff --git a/tangle/util/render_window.h b/tangle/util/render_window.h
new file mode 100644
index 0000000..4392150
--- /dev/null
+++ b/tangle/util/render_window.h
@@ -0,0 +1,95 @@
+#pragma once
+
+#include <SFML/Graphics.hpp>
+#include <SFML/Graphics/Image.hpp>
+
+#include <imgui.h>
+#include <imgui-SFML.h>
+
+#include "texture.h"
+#include "assets.h"
+
+class RenderWindow {
+private:
+sf::RenderWindow _window;
+
+sf::Sprite          _render_sprite;
+sf::Texture         _render_texture;
+cudaSurfaceObject_t _render_surface;
+sf::Rect<int>       _render_texture_view;
+
+sf::Shader _blur_shader;
+bool _blur = false;
+
+sf::Clock _ui_delta_clock;
+
+public:
+RenderWindow(std::string name):
+  _window(sf::VideoMode(800, 600), name) {
+  _render_texture.create(sf::VideoMode::getDesktopMode().width, sf::VideoMode::getDesktopMode().height);
+  _render_surface = bindTextureToCuda(_render_texture);
+  _render_sprite.setTexture(_render_texture);
+  _render_texture_view = sf::Rect<int>(0,0,_window.getSize().x,_window.getSize().y);
+  _render_sprite.setTextureRect(_render_texture_view);
+  _window.setView(sf::View(sf::Vector2f(_render_texture_view.width/2, _render_texture_view.height/2),
+                           sf::Vector2f(_window.getSize().x, _window.getSize().y)));
+  _window.setVerticalSyncEnabled(true);
+  _blur_shader.loadFromMemory(std::string(reinterpret_cast<const char*>(assets::shader::file_blur_frag)), sf::Shader::Fragment);
+  _blur_shader.setUniform("texture", sf::Shader::CurrentTexture);
+  ImGui::SFML::Init(_window);
+  ImGuiIO& io = ImGui::GetIO();
+  io.MouseDrawCursor = true;
+};
+
+bool isOpen() const {
+  return _window.isOpen();
+}
+
+void setBlur(bool state) {
+  _blur = state;
+}
+
+template <typename UI, typename MOUSE>
+void draw(UI ui, MOUSE mouse);
+
+cudaSurfaceObject_t getRenderSurface() {
+  return _render_surface;
+}
+
+sf::Rect<int> getRenderView() {
+  return _render_texture_view;
+}
+
+};
+
+template <typename UI, typename MOUSE>
+void RenderWindow::draw(UI ui, MOUSE mouse) {
+  sf::Event event;
+  while (_window.pollEvent(event)) {
+    ImGui::SFML::ProcessEvent(event);
+    if (event.type == sf::Event::Closed) {
+      _window.close();
+    }
+    if (event.type == sf::Event::Resized) {
+      _render_texture_view = sf::Rect<int>(0,0,event.size.width,event.size.height);
+      _render_sprite.setTextureRect(_render_texture_view);
+      sf::View view(sf::Vector2f(_render_texture_view.width/2, _render_texture_view.height/2),
+                    sf::Vector2f(event.size.width, event.size.height));
+      _window.setView(view);
+    }
+    if (!ImGui::GetIO().WantCaptureMouse) {
+      mouse(event);
+    }
+  }
+
+  ImGui::SFML::Update(_window, _ui_delta_clock.restart());
+  ui();
+  _window.clear();
+  if (_blur) {
+    _window.draw(_render_sprite, &_blur_shader);
+  } else {
+    _window.draw(_render_sprite);
+  }
+  ImGui::SFML::Render(_window);
+  _window.display();
+}
diff --git a/tangle/util/texture.h b/tangle/util/texture.h
new file mode 100644
index 0000000..f2e0440
--- /dev/null
+++ b/tangle/util/texture.h
@@ -0,0 +1,25 @@
+#pragma once
+
+#include <cstring>
+#include <SFML/Graphics.hpp>
+#include <cuda_gl_interop.h>
+#include <LLBM/memory.h>
+
+cudaSurfaceObject_t bindTextureToCuda(sf::Texture& texture) {
+  GLuint gl_tex_handle = texture.getNativeHandle();
+  cudaGraphicsResource* cuda_tex_handle;
+  cudaArray* buffer;
+
+  cudaGraphicsGLRegisterImage(&cuda_tex_handle, gl_tex_handle, GL_TEXTURE_2D, cudaGraphicsRegisterFlagsNone);
+  cudaGraphicsMapResources(1, &cuda_tex_handle, 0);
+  cudaGraphicsSubResourceGetMappedArray(&buffer, cuda_tex_handle, 0, 0);
+
+  cudaResourceDesc resDesc;
+  resDesc.resType = cudaResourceTypeArray;
+
+  resDesc.res.array.array = buffer;
+  cudaSurfaceObject_t cudaSurfaceObject = 0;
+  cudaCreateSurfaceObject(&cudaSurfaceObject, &resDesc);
+
+  return cudaSurfaceObject;
+}
diff --git a/tangle/util/timer.h b/tangle/util/timer.h
new file mode 100644
index 0000000..fd5d832
--- /dev/null
+++ b/tangle/util/timer.h
@@ -0,0 +1,19 @@
+#pragma once
+#include <chrono>
+
+namespace timer {
+
+std::chrono::time_point<std::chrono::steady_clock> now() {
+  return std::chrono::steady_clock::now();
+}
+
+double secondsSince(
+  std::chrono::time_point<std::chrono::steady_clock>& pit) {
+  return std::chrono::duration_cast<std::chrono::duration<double>>(now() - pit).count();
+}
+
+double mlups(std::size_t nCells, std::size_t nSteps, std::chrono::time_point<std::chrono::steady_clock>& start) {
+  return nCells * nSteps / (secondsSince(start) * 1e6);
+}
+
+}
diff --git a/tangle/util/volumetric_example.h b/tangle/util/volumetric_example.h
new file mode 100644
index 0000000..cc6a24e
--- /dev/null
+++ b/tangle/util/volumetric_example.h
@@ -0,0 +1,121 @@
+#pragma once
+#include <LLBM/volumetric.h>
+
+#include "camera.h"
+#include "texture.h"
+#include "colormap.h"
+#include "noise.h"
+#include "render_window.h"
+#include "../sampler/sampler.h"
+
+class VolumetricExample : public RenderWindow {
+private:
+std::vector<std::unique_ptr<Sampler>> _sampler;
+Sampler* _current = nullptr;
+
+Camera _camera;
+VolumetricRenderConfig _config;
+ColorPalette _palette;
+NoiseSource _noise;
+
+int _steps_per_second = 100;
+int _samples_per_second = 30;
+
+public:
+VolumetricExample(descriptor::CuboidD<3> cuboid):
+  RenderWindow("LiterateLB"),
+  _camera(make_float3(cuboid.nX/2,cuboid.nY/2,cuboid.nZ/2), cuboid.nX, M_PI/2, M_PI/2),
+  _config(cuboid),
+  _palette(_config.palette),
+  _noise(_config.noise)
+{
+  _config.canvas = this->getRenderSurface();
+  this->setBlur(_config.apply_blur);
+}
+
+template <template<typename...> class SAMPLER, typename... ARGS>
+void add(ARGS&&... args) {
+  _sampler.emplace_back(new SAMPLER(std::forward<ARGS>(args)...));
+  _current = _sampler.back().get();
+}
+
+template <typename TIMESTEP>
+void run(TIMESTEP step) {
+  sf::Clock last_sample;
+  sf::Clock last_frame;
+  std::size_t iStep = 0;
+  volatile bool simulate = true;
+
+  sf::Thread simulation([&]() {
+    while (this->isOpen()) {
+      if (last_sample.getElapsedTime().asSeconds() > 1.0 / _samples_per_second) {
+        _current->sample();
+        cudaStreamSynchronize(cudaStreamPerThread);
+        last_sample.restart();
+        if (simulate) {
+          for (unsigned i=0; i < (1.0 / _samples_per_second) * _steps_per_second; ++i) {
+            step(iStep++);
+          }
+        }
+      }
+    }
+  });
+  simulation.launch();
+
+  while (this->isOpen()) {
+    this->draw(
+      [&](){
+        ImGui::Begin("Simulation", 0, ImGuiWindowFlags_AlwaysAutoResize);
+        if (ImGui::BeginCombo("Source", _current->getName().c_str())) {
+          for (auto& option : _sampler) {
+            if (ImGui::Selectable(option->getName().c_str(), _current == option.get())) {
+              _current = option.get();
+            }
+          }
+          ImGui::EndCombo();
+        }
+        _current->interact();
+        ImGui::SliderInt("Timestep/s", &_steps_per_second, 1, 2000);
+        ImGui::SliderInt("Samples/s", &_samples_per_second, 1, 60);
+        if (simulate) {
+          simulate = !ImGui::Button("Pause");
+        } else {
+          simulate =  ImGui::Button("Continue");
+        }
+        ImGui::End();
+        ImGui::Begin("Render", 0, ImGuiWindowFlags_AlwaysAutoResize);
+        ImGui::SliderFloat("Brightness", &_config.brightness, 0.1f, 2.f);
+        ImGui::SliderFloat("Delta", &_config.delta, 0.05f, 2.f);
+        ImGui::SliderFloat("Transparency", &_config.transparency, 0.001f, 1.f);
+        _palette.interact();
+        if (ImGui::CollapsingHeader("Details")) {
+          ImGui::Checkbox("Align slices to view", &_config.align_slices_to_view);
+          ImGui::SameLine();
+          ImGui::Checkbox("Jitter", &_config.apply_noise);
+          ImGui::SameLine();
+          ImGui::Checkbox("Blur", &_config.apply_blur);
+          this->setBlur(_config.apply_blur);
+          if (_config.apply_noise) {
+            _noise.interact();
+          }
+        }
+        ImGui::End();
+      },
+      [&](sf::Event& event) {
+        _camera.handle(event);
+        _config.eye_pos = _camera.getEyePosition();
+        _config.eye_dir = _camera.getDirection();
+        _config.canvas_size = make_uint2(this->getRenderView().width, this->getRenderView().height);
+      }
+    );
+    if (last_frame.getElapsedTime().asSeconds() > 1.0 / _samples_per_second) {
+      _current->render(_config);
+      cudaStreamSynchronize(cudaStreamPerThread);
+      last_frame.restart();
+    }
+  }
+
+  simulation.wait();
+}
+
+};