Extract first public LiterateLB version

5 files changed, 345 insertions, 0 deletions

diff --git a/tangle/sampler/curl_norm.h b/tangle/sampler/curl_norm.h
new file mode 100644
index 0000000..d079a4f
--- /dev/null
+++ b/tangle/sampler/curl_norm.h
@@ -0,0 +1,77 @@
+#pragma once
+
+#include "sampler.h"
+
+#include <LLBM/kernel/collect_moments.h>
+#include <LLBM/kernel/collect_curl.h>
+
+#include <thrust/pair.h>
+#include <thrust/device_vector.h>
+#include <thrust/extrema.h>
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+class CurlNormS : public Sampler {
+private:
+Lattice<DESCRIPTOR,T,S>& _lattice;
+DeviceBuffer<bool>& _mask;
+SDF _geometry;
+
+DeviceBuffer<float> _moments_rho;
+DeviceBuffer<float> _moments_u;
+DeviceBuffer<float> _curl_norm;
+
+float _scale = 1;
+float _lower = 0;
+float _upper = 1;
+
+public:
+CurlNormS(Lattice<DESCRIPTOR,T,S>& lattice, DeviceBuffer<bool>& mask, SDF geometry):
+  Sampler("Curl norm", lattice.cuboid()),
+  _lattice(lattice),
+  _mask(mask),
+  _geometry(geometry),
+  _moments_rho(lattice.cuboid().volume),
+  _moments_u(DESCRIPTOR::d * lattice.cuboid().volume),
+  _curl_norm(lattice.cuboid().volume)
+{ }
+
+void sample() {
+  _lattice.template inspect<CollectMomentsF>(_mask, _moments_rho.device(), _moments_u.device());
+  _lattice.template inspect<CollectCurlF>(_mask, _moments_u.device(), _sample_surface, _curl_norm.device());
+}
+
+void render(VolumetricRenderConfig& config) {
+  raymarch<<<
+    dim3(config.canvas_size.x / 32 + 1, config.canvas_size.y / 32 + 1),
+    dim3(32, 32)
+  >>>(config,
+      _geometry,
+      [samples=_sample_texture, scale=_scale, lower=_lower, upper=_upper]
+      __device__ (float3 p) -> float {
+        float sample = scale * tex3D<float>(samples, p.x, p.y, p.z);
+        return sample >= lower && sample <= upper ? sample : 0;
+      },
+      [] __device__ (float x) -> float {
+        return x;
+      });
+}
+
+void scale() {
+  auto max = thrust::max_element(thrust::device_pointer_cast(_curl_norm.device()),
+                                 thrust::device_pointer_cast(_curl_norm.device() + _lattice.cuboid().volume));
+  _scale = 1 / max[0];
+}
+
+void interact() {
+  ImGui::SliderFloat("Scale", &_scale, 0.01f, 100.f);
+  ImGui::SameLine();
+  if (ImGui::Button("Auto")) {
+    scale();
+  }
+  ImGui::DragFloatRange2("Bounds", &_lower, &_upper, 0.01f, 0.f, 1.f, "%.2f", "%.2f");
+}
+
+};
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+CurlNormS(Lattice<DESCRIPTOR,T,S>&, DeviceBuffer<bool>&, SDF) -> CurlNormS<DESCRIPTOR,T,S,SDF>;
diff --git a/tangle/sampler/q_criterion.h b/tangle/sampler/q_criterion.h
new file mode 100644
index 0000000..31eeb61
--- /dev/null
+++ b/tangle/sampler/q_criterion.h
@@ -0,0 +1,90 @@
+#pragma once
+
+#include "sampler.h"
+
+#include <LLBM/kernel/collect_moments.h>
+#include <LLBM/kernel/collect_curl.h>
+#include <LLBM/kernel/collect_q_criterion.h>
+
+#include <thrust/pair.h>
+#include <thrust/device_vector.h>
+#include <thrust/extrema.h>
+
+#include <iostream>
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+class QCriterionS : public Sampler {
+private:
+Lattice<DESCRIPTOR,T,S>& _lattice;
+DeviceBuffer<bool>& _mask;
+SDF _geometry;
+
+DeviceTexture<float> _curl_buffer;
+cudaTextureObject_t _curl_texture;
+cudaSurfaceObject_t _curl_surface;
+
+DeviceBuffer<float> _moments_rho;
+DeviceBuffer<float> _moments_u;
+DeviceBuffer<float> _curl_norm;
+DeviceBuffer<float> _q;
+
+float _scale = 1;
+float _lower = 0.01;
+float _upper = 1;
+
+public:
+QCriterionS(Lattice<DESCRIPTOR,T,S>& lattice, DeviceBuffer<bool>& mask, SDF geometry):
+  Sampler("Q criterion", lattice.cuboid()),
+  _lattice(lattice),
+  _mask(mask),
+  _geometry(geometry),
+  _curl_buffer(lattice.cuboid()),
+  _curl_texture(_curl_buffer.getTexture()),
+  _curl_surface(_curl_buffer.getSurface()),
+  _moments_rho(lattice.cuboid().volume),
+  _moments_u(DESCRIPTOR::d * lattice.cuboid().volume),
+  _curl_norm(lattice.cuboid().volume),
+  _q(lattice.cuboid().volume)
+{ }
+
+void sample() {
+  _lattice.template inspect<CollectMomentsF>(_mask, _moments_rho.device(), _moments_u.device());
+  _lattice.template inspect<CollectCurlF>(_mask, _moments_u.device(), _curl_surface, _curl_norm.device());
+  _lattice.template inspect<CollectQCriterionF>(_mask, _moments_rho.device(), _moments_u.device(), _curl_norm.device(), _sample_surface, _q.device());
+}
+
+void render(VolumetricRenderConfig& config) {
+  raymarch<<<
+    dim3(config.canvas_size.x / 32 + 1, config.canvas_size.y / 32 + 1),
+    dim3(32, 32)
+  >>>(config,
+      _geometry,
+      [samples=_sample_texture, scale=_scale, lower=_lower, upper=_upper]
+      __device__ (float3 p) -> float {
+        float sample = scale * tex3D<float>(samples, p.x, p.y, p.z);
+        return (sample >= lower) * (sample <= upper) * sample;
+      },
+      [] __device__ (float x) -> float {
+        return (x > 0) * 1;
+      });
+}
+
+void scale() {
+  auto max = thrust::max_element(thrust::device_pointer_cast(_q.device()),
+                                 thrust::device_pointer_cast(_q.device() + _lattice.cuboid().volume));
+  _scale = 1 / max[0];
+}
+
+void interact() {
+  ImGui::SliderFloat("Scale", &_scale, 0.01f, 10000.f);
+  ImGui::SameLine();
+  if (ImGui::Button("Auto")) {
+    scale();
+  }
+  ImGui::DragFloatRange2("Bounds", &_lower, &_upper, 0.01f, 0.01f, 1.f, "%.2f", "%.2f");
+}
+
+};
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+QCriterionS(Lattice<DESCRIPTOR,T,S>&, DeviceBuffer<bool>&, SDF) -> QCriterionS<DESCRIPTOR,T,S,SDF>;
diff --git a/tangle/sampler/sampler.h b/tangle/sampler/sampler.h
new file mode 100644
index 0000000..ccc50b1
--- /dev/null
+++ b/tangle/sampler/sampler.h
@@ -0,0 +1,32 @@
+#pragma once
+
+#include <LLBM/base.h>
+
+class RenderWindow;
+class VolumetricRenderConfig;
+
+class Sampler {
+protected:
+const std::string _name;
+
+DeviceTexture<float> _sample_buffer;
+cudaTextureObject_t _sample_texture;
+cudaSurfaceObject_t _sample_surface;
+
+public:
+Sampler(std::string name, descriptor::CuboidD<3> cuboid):
+  _name(name),
+  _sample_buffer(cuboid),
+  _sample_texture(_sample_buffer.getTexture()),
+  _sample_surface(_sample_buffer.getSurface())
+ { }
+
+const std::string& getName() const {
+  return _name;
+}
+
+virtual void sample() = 0;
+virtual void render(VolumetricRenderConfig& config) = 0;
+virtual void interact() = 0;
+
+};
diff --git a/tangle/sampler/shear_layer.h b/tangle/sampler/shear_layer.h
new file mode 100644
index 0000000..0d42ee8
--- /dev/null
+++ b/tangle/sampler/shear_layer.h
@@ -0,0 +1,67 @@
+#pragma once
+
+#include "sampler.h"
+
+#include <LLBM/kernel/collect_moments.h>
+#include <LLBM/kernel/collect_shear_layer_normal.h>
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+class ShearLayerVisibilityS : public Sampler {
+private:
+Lattice<DESCRIPTOR,T,S>& _lattice;
+DeviceBuffer<bool>& _mask;
+SDF _geometry;
+
+DeviceBuffer<float> _moments_rho;
+DeviceBuffer<float> _moments_u;
+DeviceBuffer<float> _shear_normals;
+
+float3 _shear_layer;
+float _lower = 0;
+float _upper = 1;
+bool _center = true;
+
+public:
+ShearLayerVisibilityS(Lattice<DESCRIPTOR,T,S>& lattice, DeviceBuffer<bool>& mask, SDF geometry, float3 shear_layer):
+  Sampler("Shear layer visibility", lattice.cuboid()),
+  _lattice(lattice),
+  _mask(mask),
+  _geometry(geometry),
+  _moments_rho(lattice.cuboid().volume),
+  _moments_u(DESCRIPTOR::d * lattice.cuboid().volume),
+  _shear_normals(DESCRIPTOR::d * lattice.cuboid().volume),
+  _shear_layer(shear_layer)
+{ }
+
+void sample() {
+  _lattice.template inspect<CollectShearLayerNormalsF>(_mask, _moments_rho.device(), _moments_u.device(), _shear_normals.device());
+  _lattice.template helper<CollectShearLayerVisibilityF>(_mask, _shear_normals.device(), _shear_layer, _sample_surface);
+}
+
+void render(VolumetricRenderConfig& config) {
+  raymarch<<<
+    dim3(config.canvas_size.x / 32 + 1, config.canvas_size.y / 32 + 1),
+    dim3(32, 32)
+  >>>(config,
+      _geometry,
+      [samples=_sample_texture, lower=_lower, upper=_upper, center=_center]
+      __device__ (float3 p) -> float {
+        float sample = tex3D<float>(samples, p.x, p.y, p.z);
+        float centered = center ? 0.5 + 0.5*sample : sample;
+        return fabs(sample) >= lower && fabs(sample) <= upper ? fabs(centered) : 0;
+      },
+      [] __device__ (float x) -> float {
+        return x;
+      });
+}
+
+void interact() {
+  ImGui::InputFloat3("Normal", reinterpret_cast<float*>(&_shear_layer));
+  ImGui::Checkbox("Center", &_center);
+  ImGui::DragFloatRange2("Bounds", &_lower, &_upper, 0.01f, 0.f, 1.f, "%.2f", "%.2f");
+}
+
+};
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+ShearLayerVisibilityS(Lattice<DESCRIPTOR,T,S>&, DeviceBuffer<bool>&, SDF) -> ShearLayerVisibilityS<DESCRIPTOR,T,S,SDF>;
diff --git a/tangle/sampler/velocity_norm.h b/tangle/sampler/velocity_norm.h
new file mode 100644
index 0000000..59ca0a7
--- /dev/null
+++ b/tangle/sampler/velocity_norm.h
@@ -0,0 +1,79 @@
+#pragma once
+
+#include "sampler.h"
+
+#include <LLBM/kernel/collect_moments.h>
+#include <LLBM/kernel/collect_velocity_norm.h>
+
+#include <thrust/pair.h>
+#include <thrust/device_vector.h>
+#include <thrust/extrema.h>
+
+#include <iostream>
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+class VelocityNormS : public Sampler {
+private:
+Lattice<DESCRIPTOR,T,S>& _lattice;
+DeviceBuffer<bool>& _mask;
+SDF _geometry;
+
+DeviceBuffer<float> _moments_rho;
+DeviceBuffer<float> _moments_u;
+DeviceBuffer<float> _u_norm;
+
+float _scale = 1;
+float _lower = 0;
+float _upper = 1;
+
+public:
+VelocityNormS(Lattice<DESCRIPTOR,T,S>& lattice, DeviceBuffer<bool>& mask, SDF geometry):
+  Sampler("Velocity norm", lattice.cuboid()),
+  _lattice(lattice),
+  _mask(mask),
+  _geometry(geometry),
+  _moments_rho(lattice.cuboid().volume),
+  _moments_u(DESCRIPTOR::d * lattice.cuboid().volume),
+  _u_norm(lattice.cuboid().volume)
+{ }
+
+void sample() {
+  _lattice.template inspect<CollectMomentsF>(_mask, _moments_rho.device(), _moments_u.device());
+  _lattice.template helper<CollectVelocityNormF>(_mask, _moments_u.device(), _sample_surface, _u_norm.device());
+}
+
+void render(VolumetricRenderConfig& config) {
+  raymarch<<<
+    dim3(config.canvas_size.x / 32 + 1, config.canvas_size.y / 32 + 1),
+    dim3(32, 32)
+  >>>(config,
+      _geometry,
+      [samples=_sample_texture, scale=_scale, lower=_lower, upper=_upper]
+      __device__ (float3 p) -> float {
+        float sample = scale * tex3D<float>(samples, p.x, p.y, p.z);
+        return sample >= lower && sample <= upper ? sample : 0;
+      },
+      [] __device__ (float x) -> float {
+        return x;
+      });
+}
+
+void scale() {
+  auto max = thrust::max_element(thrust::device_pointer_cast(_u_norm.device()),
+                                 thrust::device_pointer_cast(_u_norm.device() + _lattice.cuboid().volume));
+  _scale = 1 / max[0];
+}
+
+void interact() {
+  ImGui::SliderFloat("Scale", &_scale, 0.01f, 100.f);
+  ImGui::SameLine();
+  if (ImGui::Button("Auto")) {
+    scale();
+  }
+  ImGui::DragFloatRange2("Bounds", &_lower, &_upper, 0.01f, 0.f, 1.f, "%.2f", "%.2f");
+}
+
+};
+
+template <typename DESCRIPTOR, typename T, typename S, typename SDF>
+VelocityNormS(Lattice<DESCRIPTOR,T,S>&, DeviceBuffer<bool>&, SDF) -> VelocityNormS<DESCRIPTOR,T,S,SDF>;