summaryrefslogtreecommitdiff
path: root/tangle/sampler
diff options
context:
space:
mode:
Diffstat (limited to 'tangle/sampler')
-rw-r--r--tangle/sampler/curl_norm.h77
-rw-r--r--tangle/sampler/q_criterion.h90
-rw-r--r--tangle/sampler/sampler.h32
-rw-r--r--tangle/sampler/shear_layer.h67
-rw-r--r--tangle/sampler/velocity_norm.h79
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>;