From 910c2299f4be8ec1819cddaae508d60cddcb9796 Mon Sep 17 00:00:00 2001
From: Adrian Kummerlaender
Date: Sat, 9 Nov 2019 20:44:20 +0100
Subject: Add support for SSS pattern to C++ LDC example
---
lid_driven_cavity/cpp/CMakeLists.txt | 3 +-
lid_driven_cavity/cpp/config.py | 10 +-
lid_driven_cavity/cpp/generate.py | 5 +-
lid_driven_cavity/cpp/ldc.cpp.SSS.mako | 193 +++++++++++++++++++++++++++++++++
shell.nix | 2 +-
5 files changed, 205 insertions(+), 8 deletions(-)
create mode 100644 lid_driven_cavity/cpp/ldc.cpp.SSS.mako
diff --git a/lid_driven_cavity/cpp/CMakeLists.txt b/lid_driven_cavity/cpp/CMakeLists.txt
index cc526e3..634ea58 100644
--- a/lid_driven_cavity/cpp/CMakeLists.txt
+++ b/lid_driven_cavity/cpp/CMakeLists.txt
@@ -18,8 +18,9 @@ add_custom_command(
DEPENDS
generate.py
config.py
- ldc.cpp.AA.mako
ldc.cpp.AB.mako
+ ldc.cpp.AA.mako
+ ldc.cpp.SSS.mako
)
include_directories(
diff --git a/lid_driven_cavity/cpp/config.py b/lid_driven_cavity/cpp/config.py
index 4ea30b4..0ea626b 100644
--- a/lid_driven_cavity/cpp/config.py
+++ b/lid_driven_cavity/cpp/config.py
@@ -5,11 +5,11 @@ descriptor = D2Q9
geometry = Geometry(256, 256)
tau = 0.52
precision = 'double'
-streaming = 'AA'
+streaming = 'SSS'
## 3D LDC
-#descriptor = D3Q19
+#descriptor = D3Q27
#geometry = Geometry(64, 64, 64)
-#tau = 0.52
-#precision = 'single'
-#streaming = 'AA'
+#tau = 0.6
+#precision = 'double'
+#streaming = 'SSS'
diff --git a/lid_driven_cavity/cpp/generate.py b/lid_driven_cavity/cpp/generate.py
index dae2514..973a659 100755
--- a/lid_driven_cavity/cpp/generate.py
+++ b/lid_driven_cavity/cpp/generate.py
@@ -20,13 +20,16 @@ generator = Generator(
precision = config.precision,
streaming = config.streaming,
index = 'XYZ',
- layout = 'AOS')
+ layout = 'SOA')
if args.output is None:
args.output = '.'
functions = ['collide_and_stream', 'equilibrilize', 'collect_moments', 'momenta_boundary']
+if config.streaming == 'SSS':
+ functions += ['update_sss_control_structure']
+
with open('%s/kernel.h' % args.output, 'w') as kernel:
kernel.write(generator.kernel(config.geometry, functions))
diff --git a/lid_driven_cavity/cpp/ldc.cpp.SSS.mako b/lid_driven_cavity/cpp/ldc.cpp.SSS.mako
new file mode 100644
index 0000000..ed2f938
--- /dev/null
+++ b/lid_driven_cavity/cpp/ldc.cpp.SSS.mako
@@ -0,0 +1,193 @@
+#include <cstdint>
+#include <memory>
+#include <vector>
+#include <chrono>
+#include <iostream>
+#include <fstream>
+
+#include "kernel.h"
+
+void collect_moments_to_vtk(const std::string& path, ${float_type}** f) {
+ std::ofstream fout;
+ fout.open(path.c_str());
+
+ fout << "# vtk DataFile Version 3.0\n";
+ fout << "lbm_output\n";
+ fout << "ASCII\n";
+ fout << "DATASET RECTILINEAR_GRID\n";
+% if descriptor.d == 2:
+ fout << "DIMENSIONS " << ${geometry.size_x-2} << " " << ${geometry.size_y-2} << " 1" << "\n";
+% else:
+ fout << "DIMENSIONS " << ${geometry.size_x-2} << " " << ${geometry.size_y-2} << " " << ${geometry.size_z-2} << "\n";
+% endif
+
+ fout << "X_COORDINATES " << ${geometry.size_x-2} << " float\n";
+ for( std::size_t x = 1; x < ${geometry.size_x-1}; ++x ) {
+ fout << x << " ";
+ }
+
+ fout << "\nY_COORDINATES " << ${geometry.size_y-2} << " float\n";
+ for( std::size_t y = 1; y < ${geometry.size_y-1}; ++y ) {
+ fout << y << " ";
+ }
+
+% if descriptor.d == 2:
+ fout << "\nZ_COORDINATES " << 1 << " float\n";
+ fout << 0 << "\n";
+ fout << "POINT_DATA " << ${(geometry.size_x-2) * (geometry.size_y-2)} << "\n";
+% else:
+ fout << "\nZ_COORDINATES " << ${geometry.size_z-2} << " float\n";
+ for( std::size_t z = 1; z < ${geometry.size_z-1}; ++z ) {
+ fout << z << " ";
+ }
+ fout << "\nPOINT_DATA " << ${(geometry.size_x-2) * (geometry.size_y-2) * (geometry.size_z-2)} << "\n";
+% endif
+
+ ${float_type} rho;
+ ${float_type} u[${descriptor.d}];
+
+ fout << "VECTORS velocity float\n";
+% if descriptor.d == 2:
+ for ( std::size_t y = 1; y < ${geometry.size_y-1}; ++y ) {
+ for ( std::size_t x = 1; x < ${geometry.size_x-1}; ++x ) {
+ collect_moments(f, x*${geometry.size_y}+y, rho, u);
+ fout << u[0] << " " << u[1] << " 0\n";
+ }
+ }
+% else:
+ for ( std::size_t z = 1; z < ${geometry.size_z-1}; ++z ) {
+ for ( std::size_t y = 1; y < ${geometry.size_y-1}; ++y ) {
+ for ( std::size_t x = 1; x < ${geometry.size_x-1}; ++x ) {
+ collect_moments(f, x*${geometry.size_y*geometry.size_z}+y*${geometry.size_z}+z, rho, u);
+ fout << u[0] << " " << u[1] << " " << u[2] << "\n";
+ }
+ }
+ }
+% endif
+
+ fout << "SCALARS density float 1\n";
+ fout << "LOOKUP_TABLE default\n";
+% if descriptor.d == 2:
+ for ( std::size_t y = 1; y < ${geometry.size_y-1}; ++y ) {
+ for ( std::size_t x = 1; x < ${geometry.size_x-1}; ++x ) {
+ collect_moments(f, x*${geometry.size_y}+y, rho, u);
+ fout << rho << "\n";
+ }
+ }
+% else:
+ for ( std::size_t z = 1; z < ${geometry.size_z-1}; ++z ) {
+ for ( std::size_t y = 1; y < ${geometry.size_y-1}; ++y ) {
+ for ( std::size_t x = 1; x < ${geometry.size_x-1}; ++x ) {
+ collect_moments(f, x*${geometry.size_y*geometry.size_z}+y*${geometry.size_z}+z, rho, u);
+ fout << rho << "\n";
+ }
+ }
+ }
+% endif
+
+ fout.close();
+}
+
+void simulate(std::size_t nStep)
+{
+<%
+padding = (max(geometry.size_x,geometry.size_y,geometry.size_z)+1)**(descriptor.d-1)
+%>
+ auto f_aa = std::make_unique<${float_type}[]>(${(geometry.volume+2*padding)*descriptor.q});
+
+ ${float_type}* f[${descriptor.q}];
+
+% for i, c_i in enumerate(descriptor.c):
+ f[${i}] = f_aa.get() + ${layout.pop_offset(i) + padding};
+% endfor
+
+ std::vector<std::size_t> ghost;
+ std::vector<std::size_t> bulk;
+ std::vector<std::size_t> lid_bc;
+ std::vector<std::size_t> box_bc;
+
+ for (int iX = 0; iX < ${geometry.size_x}; ++iX) {
+ for (int iY = 0; iY < ${geometry.size_y}; ++iY) {
+% if descriptor.d == 2:
+ const std::size_t iCell = iX*${geometry.size_y} + iY;
+ if (iX == 0 || iY == 0 || iX == ${geometry.size_x-1} || iY == ${geometry.size_y-1}) {
+ ghost.emplace_back(iCell);
+ } else if (iY == ${geometry.size_y-2}) {
+ lid_bc.emplace_back(iCell);
+ } else if (iX == 1 || iX == ${geometry.size_x-2} || iY == 1) {
+ box_bc.emplace_back(iCell);
+ } else {
+ bulk.emplace_back(iCell);
+ }
+% elif descriptor.d == 3:
+ for (int iZ = 0; iZ < ${geometry.size_z}; ++iZ) {
+ const std::size_t iCell = iX*${geometry.size_y*geometry.size_z} + iY*${geometry.size_z} + iZ;
+ if ( iX == 0 || iY == 0 || iZ == 0
+ || iX == ${geometry.size_x-1}
+ || iY == ${geometry.size_y-1}
+ || iZ == ${geometry.size_z-1}) {
+ ghost.emplace_back(iCell);
+ } else if (iZ == ${geometry.size_z-2}) {
+ lid_bc.emplace_back(iCell);
+ } else if ( iX == 1 || iX == ${geometry.size_x-2}
+ || iY == 1 || iY == ${geometry.size_y-2}
+ || iZ == 1) {
+ box_bc.emplace_back(iCell);
+ } else {
+ bulk.emplace_back(iCell);
+ }
+ }
+% endif
+ }
+ }
+
+ std::cout << "#ghost : " << ghost.size() << std::endl;
+ std::cout << "#bulk : " << bulk.size() << std::endl;
+ std::cout << "#lid : " << lid_bc.size() << std::endl;
+ std::cout << "#wall : " << box_bc.size() << std::endl;
+ std::cout << std::endl;
+
+ #pragma omp parallel for
+ for (std::size_t iCell = 0; iCell < ${geometry.volume}; ++iCell) {
+ equilibrilize(f, iCell);
+ }
+
+ auto start = std::chrono::high_resolution_clock::now();
+
+ for (std::size_t iStep = 1; iStep <= nStep; ++iStep) {
+ #pragma omp parallel for
+ for (std::size_t i = 0; i < ghost.size(); ++i) {
+ equilibrilize(f, ghost[i]);
+ }
+ #pragma omp parallel for
+ for (std::size_t i = 0; i < bulk.size(); ++i) {
+ collide_and_stream(f, bulk[i]);
+ }
+ ${float_type} u[${descriptor.d}] { 0. };
+ #pragma omp parallel for
+ for (std::size_t i = 0; i < box_bc.size(); ++i) {
+ velocity_momenta_boundary(f, box_bc[i], u);
+ }
+ u[0] = 0.05;
+ #pragma omp parallel for
+ for (std::size_t i = 0; i < lid_bc.size(); ++i) {
+ velocity_momenta_boundary(f, lid_bc[i], u);
+ }
+
+ if (iStep % 1000 == 0) {
+ auto duration = std::chrono::duration_cast<std::chrono::duration<double>>(
+ std::chrono::high_resolution_clock::now() - start);
+ std::cout << "iStep = " << iStep << "; ~" << 1000*${geometry.volume}/(1e6*duration.count()) << " MLUPS" << std::endl;
+
+ collect_moments_to_vtk("result/ldc_" + std::to_string(iStep) + ".vtk", f);
+
+ start = std::chrono::high_resolution_clock::now();
+ }
+
+ update_sss_control_structure(f);
+ }
+}
+
+int main() {
+ simulate(20000);
+}
diff --git a/shell.nix b/shell.nix
index e3e52d4..45805fb 100644
--- a/shell.nix
+++ b/shell.nix
@@ -27,7 +27,7 @@ pkgs.stdenvNoCC.mkDerivation rec {
src = builtins.fetchGit {
url = "https://code.kummerlaender.eu/boltzgen/";
- rev = "286e243a171c8bcdfc91b5b6dcdd937ac95b0b7b";
+ rev = "27ce855378a80dff680c2989800af1f4e69975fe";
};
propagatedBuildInputs = with pkgs.python37Packages; [
--
cgit v1.2.3