summaryrefslogtreecommitdiff
path: root/tangle/channel.cu
diff options
context:
space:
mode:
authorAdrian Kummerlaender2021-09-19 21:02:52 +0200
committerAdrian Kummerlaender2021-09-19 21:02:52 +0200
commit186bfae321f9d34b702a7c31b420eed4d10e7a50 (patch)
tree4e8e6a691193dd0ce9555d2e29f4414336d17cda /tangle/channel.cu
parent32dd41a728ce10113032e20955ba08f8de449857 (diff)
downloadLiterateLB-186bfae321f9d34b702a7c31b420eed4d10e7a50.tar
LiterateLB-186bfae321f9d34b702a7c31b420eed4d10e7a50.tar.gz
LiterateLB-186bfae321f9d34b702a7c31b420eed4d10e7a50.tar.bz2
LiterateLB-186bfae321f9d34b702a7c31b420eed4d10e7a50.tar.lz
LiterateLB-186bfae321f9d34b702a7c31b420eed4d10e7a50.tar.xz
LiterateLB-186bfae321f9d34b702a7c31b420eed4d10e7a50.tar.zst
LiterateLB-186bfae321f9d34b702a7c31b420eed4d10e7a50.zip
Small changes to channel example
Diffstat (limited to 'tangle/channel.cu')
-rw-r--r--tangle/channel.cu85
1 files changed, 85 insertions, 0 deletions
diff --git a/tangle/channel.cu b/tangle/channel.cu
new file mode 100644
index 0000000..6a0f072
--- /dev/null
+++ b/tangle/channel.cu
@@ -0,0 +1,85 @@
+#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 "util/volumetric_example.h"
+#include "sampler/velocity_norm.h"
+#include "sampler/curl_norm.h"
+#include "sampler/q_criterion.h"
+
+using T = float;
+using DESCRIPTOR = descriptor::D3Q19;
+
+int main() {
+if (cuda::device::count() == 0) {
+ std::cerr << "No CUDA devices on this system" << std::endl;
+ return -1;
+}
+auto current = cuda::device::current::get();
+
+const descriptor::Cuboid<DESCRIPTOR> cuboid(500, 100, 100);
+Lattice<DESCRIPTOR,T> lattice(cuboid);
+
+CellMaterials<DESCRIPTOR> materials(cuboid, [&cuboid](uint3 p) -> int {
+ if (p.z == 0 || p.z == cuboid.nZ-1) {
+ return 2; // boundary cell
+ } else if (p.y == 0 || p.y == cuboid.nY-1) {
+ return 3; // boundary cell
+ } else if (p.x == 0) {
+ return 4; // inflow cell
+ } else if (p.x == cuboid.nX-1) {
+ return 5; // outflow cell
+ } else {
+ return 1; // bulk
+ }
+});
+
+for (std::size_t iX=0; iX < cuboid.nX; ++iX) {
+ materials.set(gid(cuboid, iX, 0, 0), 6);
+ materials.set(gid(cuboid, iX, cuboid.nY-1, 0), 6);
+ materials.set(gid(cuboid, iX, 0, cuboid.nZ-1), 6);
+ materials.set(gid(cuboid, iX, cuboid.nY-1, cuboid.nZ-1), 6);
+}
+
+auto obstacle = [cuboid] __host__ __device__ (float3 p) -> float {
+ p -= make_float3(cuboid.nX/5, cuboid.nY/2, cuboid.nZ/2);
+ float3 q = sdf::twisted(p, 0.01);
+ return sdf::sphere(p, cuboid.nY/3.5) + sin(0.2*q.x)*sin(0.2*q.y)*sin(0.2*q.z);
+};
+
+materials.sdf(obstacle, 0);
+SignedDistanceBoundary bouzidi(lattice, materials, obstacle, 1, 0);
+
+auto bulk_mask = materials.mask_of_material(1);
+auto wall_mask_z = materials.mask_of_material(2);
+auto wall_mask_y = materials.mask_of_material(3);
+auto inflow_mask = materials.mask_of_material(4);
+auto outflow_mask = materials.mask_of_material(5);
+auto edge_mask = materials.mask_of_material(6);
+
+cuda::synchronize(current);
+
+VolumetricExample renderer(cuboid);
+renderer.add<QCriterionS>(lattice, bulk_mask, obstacle);
+renderer.add<CurlNormS>(lattice, bulk_mask, obstacle);
+renderer.add<VelocityNormS>(lattice, bulk_mask, obstacle);
+renderer.run([&](std::size_t iStep) {
+ const float tau = 0.501;
+ const float smagorinsky = 0.1;
+ const float inflow = 0.04;
+
+ lattice.apply(Operator(SmagorinskyBgkCollideO(), bulk_mask, tau, smagorinsky),
+ Operator(BounceBackFreeSlipO(), wall_mask_z, WallNormal<0,0,1>()),
+ Operator(BounceBackFreeSlipO(), wall_mask_y, WallNormal<0,1,0>()),
+ Operator(EquilibriumVelocityWallO(), inflow_mask, std::min(iStep*1e-4, 1.0)*inflow, WallNormal<1,0,0>()),
+ Operator(EquilibriumDensityWallO(), outflow_mask, 1, WallNormal<-1,0,0>()),
+ Operator(BounceBackO(), edge_mask));
+ lattice.apply<BouzidiO>(bouzidi.getCount(), bouzidi.getConfig());
+
+ lattice.stream();
+});
+}