summaryrefslogtreecommitdiff
path: root/tangle/magnus.cu
diff options
context:
space:
mode:
authorAdrian Kummerlaender2021-05-17 00:15:33 +0200
committerAdrian Kummerlaender2021-05-17 00:15:33 +0200
commit4ec94c97879aafef15f7663135745e4ba61e62cf (patch)
tree322ae3f003892513f529842ff0b3fd100573b680 /tangle/magnus.cu
downloadLiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar
LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.gz
LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.bz2
LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.lz
LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.xz
LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.zst
LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.zip
Extract first public LiterateLB version
Diffstat (limited to 'tangle/magnus.cu')
-rw-r--r--tangle/magnus.cu116
1 files changed, 116 insertions, 0 deletions
diff --git a/tangle/magnus.cu b/tangle/magnus.cu
new file mode 100644
index 0000000..aa31ba1
--- /dev/null
+++ b/tangle/magnus.cu
@@ -0,0 +1,116 @@
+#include <LLBM/base.h>
+#include <LLBM/bulk.h>
+#include <LLBM/boundary.h>
+
+#include "util/render_window.h"
+#include "util/texture.h"
+#include "util/colormap.h"
+
+#include <LLBM/kernel/collect_moments.h>
+#include <LLBM/kernel/collect_velocity_norm.h>
+
+using T = float;
+using DESCRIPTOR = descriptor::D2Q9;
+
+int main() {
+cudaSetDevice(0);
+
+const descriptor::Cuboid<DESCRIPTOR> cuboid(1200, 500);
+Lattice<DESCRIPTOR,T> lattice(cuboid);
+
+const float tau = 0.54;
+const float u_inflow = 0.02;
+const float u_rotate = 0.08;
+
+CellMaterials<DESCRIPTOR> materials(cuboid, [&cuboid](uint2 p) -> int {
+ if (p.x == 0) {
+ return 3; // inflow
+ } else if (p.x == cuboid.nX-1) {
+ return 4; // outflow
+ } else if (p.y == 0 || p.y == cuboid.nY-1) {
+ return 2; // wall
+ } else {
+ return 1; // bulk
+ }
+});
+
+materials.set(gid(cuboid, 0,0), 2);
+materials.set(gid(cuboid, 0,cuboid.nY-1), 2);
+materials.set(gid(cuboid, cuboid.nX-1,0), 5);
+materials.set(gid(cuboid, cuboid.nX-1,cuboid.nY-1), 5);
+
+auto cylinder = [cuboid] __host__ __device__ (float2 p) -> float {
+ float2 q = p - make_float2(cuboid.nX/6, 3*cuboid.nY/4);
+ float2 r = p - make_float2(cuboid.nX/6, 1*cuboid.nY/4);
+ return sdf::add(sdf::sphere(q, cuboid.nY/18),
+ sdf::sphere(r, cuboid.nY/18));
+ };
+
+materials.sdf(cylinder, 0);
+SignedDistanceBoundary bouzidi(lattice, materials, cylinder, 1, 0);
+
+bouzidi.setVelocity([cuboid,u_rotate](float2 p) -> float2 {
+ float2 q = p - make_float2(cuboid.nX/6, 3*cuboid.nY/4);
+ if (length(q) < 1.1*cuboid.nY/18) {
+ return u_rotate * normalize(make_float2(-q.y, q.x));
+ } else {
+ return make_float2(0);
+ }
+});
+
+auto bulk_mask = materials.mask_of_material(1);
+auto wall_mask = materials.mask_of_material(2);
+auto inflow_mask = materials.mask_of_material(3);
+auto outflow_mask = materials.mask_of_material(4);
+auto edge_mask = materials.mask_of_material(5);
+
+lattice.apply(Operator(InitializeO(), bulk_mask),
+ Operator(InitializeO(), wall_mask),
+ Operator(InitializeO(), inflow_mask),
+ Operator(InitializeO(), outflow_mask),
+ Operator(InitializeO(), edge_mask));
+cudaDeviceSynchronize();
+
+RenderWindow window("Magnus");
+cudaSurfaceObject_t colormap;
+ColorPalette palette(colormap);
+auto slice = [cuboid] __device__ (int iX, int iY) -> std::size_t {
+ return descriptor::gid(cuboid,iX,cuboid.nY-1-iY);
+ };
+DeviceBuffer<T> moments_rho(cuboid.volume);
+DeviceBuffer<T> moments_u(2*cuboid.volume);
+T* u = moments_u.device();
+std::size_t iStep = 0;
+
+while (window.isOpen()) {
+ lattice.apply(Operator(BgkCollideO(), bulk_mask, tau),
+ Operator(BounceBackFreeSlipO(), wall_mask, WallNormal<0,1>()),
+ Operator(EquilibriumVelocityWallO(), inflow_mask, std::min(iStep*1e-5, 1.)*u_inflow, WallNormal<1,0>()),
+ Operator(EquilibriumDensityWallO(), outflow_mask, 1., WallNormal<-1,0>()),
+ Operator(BounceBackO(), edge_mask));
+ lattice.apply<BouzidiO>(bouzidi.getCount(), bouzidi.getConfig());
+ lattice.stream();
+ if (iStep % 100 == 0) {
+ cudaDeviceSynchronize();
+ lattice.inspect<CollectMomentsF>(bulk_mask, moments_rho.device(), moments_u.device());
+ renderSliceViewToTexture<<<
+ dim3(cuboid.nX / 32 + 1, cuboid.nY / 32 + 1),
+ dim3(32,32)
+ >>>(cuboid.nX, cuboid.nY,
+ slice,
+ [u,u_rotate] __device__ (std::size_t gid) -> float {
+ return length(make_float2(u[2*gid+0], u[2*gid+1])) / u_rotate;
+ },
+ [colormap] __device__ (float x) -> float3 {
+ return colorFromTexture(colormap, clamp(x, 0.f, 1.f));
+ },
+ window.getRenderSurface());
+ window.draw([&]() {
+ ImGui::Begin("Render");
+ palette.interact();
+ ImGui::End();
+ }, [](sf::Event&) { });
+ }
+ ++iStep;
+}
+}