From 3e7f672b35f4a2196be8f27d35ef527eec79179d Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Wed, 11 Feb 2026 17:28:23 +0100 Subject: Too lazy to ammend last commit. Moving to particle based HLBM and others. Also implemented rotational movement properly now --- default.nix | 2 +- examples/poiseulle_particles_2d_gpu/sim.cpp | 83 ++++++++++++++++++++--------- shell.nix | 18 +++++++ 3 files changed, 78 insertions(+), 25 deletions(-) create mode 100644 shell.nix diff --git a/default.nix b/default.nix index 7c19841..706c48f 100644 --- a/default.nix +++ b/default.nix @@ -68,7 +68,7 @@ let src = builtins.fetchurl { url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz"; - sha256 = "e91c18fef798dd7b3afbd1615c2e320b90d74aa2d7ef726801a76e3f7f77ae81"; + sha256 = "sha256:10dffxzkyvm705l75vypl95dg40b68p5qqfizcx7ppcqyzz1h779"; }; phases = [ "unpackPhase" "installPhase" ]; diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index 0ad1b61..691ec93 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -1,5 +1,6 @@ #include #include +#include #include #include @@ -9,8 +10,9 @@ namespace kel { namespace lbm { -constexpr uint64_t dim_x = 1024u; -constexpr uint64_t dim_y = 512u; +constexpr uint64_t dim_x = 1024ul; +constexpr uint64_t dim_y = 512ul; +constexpr uint64_t particle_size = 128ul; namespace sch { using namespace saw::schema; @@ -20,11 +22,16 @@ using InfoChunk = Chunk; template using DfChunk = Chunk, 1u, dim_x, dim_y>; +template +using ScalarChunk = Chunk, 0u, dim_x, dim_y>; + template using ChunkStruct = Struct< Member, Member, "dfs">, - Member, "dfs_old"> + Member, "dfs_old">, + Member, "particle_N">, + Member, "particle_D"> >; template @@ -36,17 +43,21 @@ using RhoChunk = Chunk, 0u, dim_x, dim_y>; template using MacroStruct = Struct< Member, "velocity">, - Member, "density"> + Member, "density">, + Member, "porosity"> >; -// template -// using Particles = Particle< +//template +//using ParticleArray = Array< +// Particle +//>; } template saw::error_or setup_initial_conditions( saw::data>& fields, - saw::data>& macros + saw::data>& macros, + saw::data, particle_size>>& particles ){ auto& info_f = fields.template get<"info">(); // Set everything as walls @@ -122,6 +133,10 @@ saw::error_or setup_initial_conditions( {{1u,1u}} ); + for(saw::data i{0u}; i < saw::data{particle_size}; ++i){ + auto& part = particles.at(i); + } + return saw::make_void(); } @@ -138,7 +153,7 @@ saw::error_or step( // auto coll_ev = q.submit([&](acpp::sycl::handler& h){ // Need nicer things to handle the flow. I see improvement here - component> collision{0.6}; + component> collision{0.6}; component> bb; saw::data> rho_b; @@ -232,8 +247,17 @@ saw::error_or lbm_main(int argc, char** argv){ return std::move(eo_lbm_dir.get_error()); } auto& lbm_dir = eo_lbm_dir.get_value(); + auto out_dir = lbm_dir / "poiseulle_particles_2d_gpu"; + { + std::error_code ec; + std::filesystem::create_directories(out_dir,ec); + if(ec != std::errc{}){ + return saw::make_error("Could not create output directory"); + } + } + converter conv { // delta_x {{1.0}}, @@ -244,6 +268,7 @@ saw::error_or lbm_main(int argc, char** argv){ // saw::data> meta{{dim_x,dim_y}}; auto lbm_data_ptr = saw::heap>>(); auto lbm_macro_data_ptr = saw::heap>>(); + auto lbm_particle_data_ptr = saw::heap, particle_size>>>(); auto eo_aio = saw::setup_async_io(); if(eo_aio.is_error()){ @@ -267,7 +292,7 @@ saw::error_or lbm_main(int argc, char** argv){ sycl_q.wait(); { - auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr); + auto eov = setup_initial_conditions(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr); if(eov.is_error()){ return eov; } @@ -275,7 +300,9 @@ saw::error_or lbm_main(int argc, char** argv){ saw::data, encode::Sycl> lbm_sycl_data{sycl_q}; saw::data, encode::Sycl> lbm_sycl_macro_data{sycl_q}; + saw::data, particle_size>, encode::Sycl> lbm_sycl_particle_data{sycl_q}; sycl_q.wait(); + auto lsd_view = make_chunk_struct_view(lbm_sycl_data); auto lsdm_view = make_chunk_struct_view(lbm_sycl_macro_data); { @@ -290,33 +317,41 @@ saw::error_or lbm_main(int argc, char** argv){ return eov; } } + { + auto eov = dev.copy_to_device(*lbm_particle_data_ptr,lbm_sycl_particle_data); + if(eov.is_error()){ + return eov; + } + } sycl_q.wait(); - saw::data time_steps{4096ul}; + saw::data time_steps{particle_size}; for(saw::data i{0u}; i < time_steps and krun; ++i){ { - std::string file_name = "tmp/t_"; - file_name += std::to_string(i.get()); - file_name += ".vtk"; - { - auto eov = write_vtk_file(file_name, *lbm_macro_data_ptr); + std::string file_name = "t_"; + file_name += std::to_string(i.get()); + file_name += ".vtk"; + auto eov = write_vtk_file(out_dir/file_name, *lbm_macro_data_ptr); if(eov.is_error()){ return eov; } } - { - auto eov = saw::easy::encode_and_write_file,saw::encode::KelSimple>(file_name, *lbm_macro_data_ptr); + /*{ + std::string file_name = "p_"; + file_name += std::to_string(i.get()); + file_name += ".json"; + auto eov = saw::easy::encode_and_write_file,particle_size>,saw::encode::Json>(out_dir, *lbm_particle_data_ptr); if(eov.is_error()){ return eov; } - } + }*/ } { - std::string file_name = "tmp/df_"; + std::string file_name = "df_"; file_name += std::to_string(i.get()); file_name += ".vtk"; - auto eov = write_vtk_file(file_name, *lbm_data_ptr); + auto eov = write_vtk_file(out_dir/file_name, *lbm_data_ptr); if(eov.is_error()){ return eov; } @@ -348,19 +383,19 @@ saw::error_or lbm_main(int argc, char** argv){ } sycl_q.wait(); { - std::string file_name = "tmp/t_"; + std::string file_name = "t_"; file_name += std::to_string(time_steps.get()); file_name += ".vtk"; - auto eov = write_vtk_file(file_name, *lbm_macro_data_ptr); + auto eov = write_vtk_file(out_dir/file_name, *lbm_macro_data_ptr); if(eov.is_error()){ return eov; } } { - std::string file_name = "tmp/df_"; + std::string file_name = "df_"; file_name += std::to_string(time_steps.get()); file_name += ".vtk"; - auto eov = write_vtk_file(file_name, *lbm_data_ptr); + auto eov = write_vtk_file(out_dir/file_name, *lbm_data_ptr); if(eov.is_error()){ return eov; } diff --git a/shell.nix b/shell.nix new file mode 100644 index 0000000..0c22e94 --- /dev/null +++ b/shell.nix @@ -0,0 +1,18 @@ +{ pkgs ? import {} +, ... +}: + +pkgs.mkShell { + buildInputs = [ + pkgs.scons + pkgs.clang + pkgs.clang-tools + ]; + + shellHook = '' + if [ ! -f compile_commands.json ]; then + printf "Generating compile_commands.json...\n" + scons cdb + fi + ''; +} -- cgit v1.2.3