From 1dda51f5a95d7dc6e8bcf617f53cb763753db0a7 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Sat, 11 Jul 2026 17:17:22 +0200 Subject: moving test --- .../moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp | 22 ---------------------- 1 file changed, 22 deletions(-) diff --git a/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp index ce93d3e..daa397a 100644 --- a/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp +++ b/examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp @@ -144,28 +144,6 @@ saw::error_or setup_initial_conditions( df_f.get_dims(), {{1u,1u}} ); - /* Technically the particle group should handle this case - iterator::apply( - [&](auto& index){ - saw::data> middle, ind_vec; - middle.at({{0u}}) = dim_x * 0.25; - middle.at({{1u}}) = dim_y * 0.5; - - ind_vec.at({{0u}}) = index.at({{0u}}).template cast_to(); - ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to(); - - auto dist = middle - ind_vec; - auto dist_2 = saw::math::dot(dist,dist); - if(dist_2.at({}).get() < dim_y*dim_y*0.01){ - porous_f.at(index).at({}) = 0.0; - } - }, - {},// 0-index - df_f.get_dims() - ); - */ - - return saw::make_void(); } -- cgit v1.2.3 From 5c9e64112e1f710e99790db4814eb2f0031dec34 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Sun, 12 Jul 2026 22:38:58 +0200 Subject: Adding poc ideas --- poc/mixed_precision/sim.cpp | 24 ++++++++++++++++++++++++ poc/triangle_stream/sim.cpp | 24 ++++++++++++++++++++++++ 2 files changed, 48 insertions(+) create mode 100644 poc/mixed_precision/sim.cpp create mode 100644 poc/triangle_stream/sim.cpp diff --git a/poc/mixed_precision/sim.cpp b/poc/mixed_precision/sim.cpp new file mode 100644 index 0000000..c96d7d0 --- /dev/null +++ b/poc/mixed_precision/sim.cpp @@ -0,0 +1,24 @@ +void init(){ +} + +saw::error_or kmain(int argc, char** argv){ + using namespace kel; + + return saw::make_void(); +} + +int main(int argc, char** argv){ + auto eov = kmain(argc,argv); + if(eov.is_error()){ + auto& err = eov.get_error(); + + std::cerr<<"[Error]: "< kmain(int argc, char** argv){ + using namespace kel; + + return saw::make_void(); +} + +int main(int argc, char** argv){ + auto eov = kmain(argc,argv); + if(eov.is_error()){ + auto& err = eov.get_error(); + + std::cerr<<"[Error]: "< Date: Mon, 13 Jul 2026 15:04:22 +0200 Subject: poc work --- poc/triangle_stream/.nix/derivation.nix | 41 +++++++++++++++++ poc/triangle_stream/SConscript | 34 ++++++++++++++ poc/triangle_stream/SConstruct | 81 +++++++++++++++++++++++++++++++++ poc/triangle_stream/sim.cpp | 60 +++++++++++++++++++++++- 4 files changed, 214 insertions(+), 2 deletions(-) create mode 100644 poc/triangle_stream/.nix/derivation.nix create mode 100644 poc/triangle_stream/SConscript create mode 100644 poc/triangle_stream/SConstruct diff --git a/poc/triangle_stream/.nix/derivation.nix b/poc/triangle_stream/.nix/derivation.nix new file mode 100644 index 0000000..c86fc0b --- /dev/null +++ b/poc/triangle_stream/.nix/derivation.nix @@ -0,0 +1,41 @@ +{ lib +, stdenv +, scons +, clang-tools +, forstio +, python3 +, pname +, version +, adaptive-cpp +, kel +}: + +stdenv.mkDerivation { + pname = pname + "-poc-" + "triangle_stream"; + inherit version; + src = ./..; + + nativeBuildInputs = [ + scons + clang-tools + python3 + ]; + + buildInputs = [ + forstio.core + forstio.async + forstio.codec + forstio.codec-unit + forstio.io + forstio.remote + forstio.remote-filesystem + forstio.codec-json + adaptive-cpp + kel.lbm.core + kel.lbm.sycl + ]; + + preferLocalBuild = true; + + outputs = [ "out" "dev" ]; +} diff --git a/poc/triangle_stream/SConscript b/poc/triangle_stream/SConscript new file mode 100644 index 0000000..992fc60 --- /dev/null +++ b/poc/triangle_stream/SConscript @@ -0,0 +1,34 @@ +#!/bin/false + +import os +import os.path +import glob + + +Import('env') + +dir_path = Dir('.').abspath + +# Environment for base library +examples_env = env.Clone(); +examples_env['CXX'] = 'syclcc-clang'; +examples_env['CXXFLAGS'] += ['-O3']; + +examples_env.sources = sorted(glob.glob(dir_path + "/*.cpp")) +examples_env.headers = sorted(glob.glob(dir_path + "/*.hpp")) + +env.sources += examples_env.sources; +env.headers += examples_env.headers; + +# Cavity2D +examples_objects = []; +examples_env.add_source_files(examples_objects, ['sim.cpp'], shared=False); +examples_env.poiseulle_2d_gpu = examples_env.Program('#bin/kel_lbm_triangle_stream', [examples_objects]); + +# Set Alias +env.examples = [ + examples_env.poiseulle_2d_gpu +]; +env.Alias('examples', env.examples); +env.targets += ['examples']; +env.Install('$prefix/bin/', env.examples); diff --git a/poc/triangle_stream/SConstruct b/poc/triangle_stream/SConstruct new file mode 100644 index 0000000..0611b67 --- /dev/null +++ b/poc/triangle_stream/SConstruct @@ -0,0 +1,81 @@ +#!/usr/bin/env python3 + +import sys +import os +import os.path +import glob +import re + + +if sys.version_info < (3,): + def isbasestring(s): + return isinstance(s,basestring) +else: + def isbasestring(s): + return isinstance(s, (str,bytes)) + +def add_kel_source_files(self, sources, filetype, lib_env=None, shared=False, target_post=""): + + if isbasestring(filetype): + dir_path = self.Dir('.').abspath + filetype = sorted(glob.glob(dir_path+"/"+filetype)) + + for path in filetype: + target_name = re.sub( r'(.*?)(\.cpp|\.c\+\+)', r'\1' + target_post, path ) + if shared: + target_name+='.os' + sources.append( self.SharedObject( target=target_name, source=path ) ) + else: + target_name+='.o' + sources.append( self.StaticObject( target=target_name, source=path ) ) + pass + +def isAbsolutePath(key, dirname, env): + assert os.path.isabs(dirname), "%r must have absolute path syntax" % (key,) + +env_vars = Variables( + args=ARGUMENTS +) + +env_vars.Add('prefix', + help='Installation target location of build results and headers', + default='/usr/local/', + validator=isAbsolutePath +) + +env_vars.Add('build_examples', + help='If examples should be built', + default="true" +) + +env=Environment(ENV=os.environ, variables=env_vars, CPPPATH=[], + CPPDEFINES=['SAW_UNIX'], + CXXFLAGS=[ + '-std=c++20', + '-g', + '-Wall', + '-Wextra' + ], + LIBS=[ + 'forstio-core', + 'forstio-async', + 'forstio-io' + ] +); +env.__class__.add_source_files = add_kel_source_files +env.Tool('compilation_db'); +env.cdb = env.CompilationDatabase('compile_commands.json'); + +env.objects = []; +env.sources = []; +env.headers = []; +env.targets = []; + +Export('env') +SConscript('SConscript') + +env.Alias('cdb', env.cdb); +env.Alias('all', [env.targets]); +env.Default('all'); + +env.Alias('install', '$prefix') diff --git a/poc/triangle_stream/sim.cpp b/poc/triangle_stream/sim.cpp index c96d7d0..36986e4 100644 --- a/poc/triangle_stream/sim.cpp +++ b/poc/triangle_stream/sim.cpp @@ -1,14 +1,70 @@ +#include + +namespace kel { +namespace lbm { + +constexpr uint64_t dim_y = 256ul; +constexpr uint64_t dim_x = dim_y * 20ul; + +namespace sch { +using namespace saw::schema; + +using InfoChunk = Chunk; + +template +using DfChunk = Chunk, 1u, dim_x, dim_y>; + +template +using ScalarChunk = Chunk, 0u, dim_x, dim_y>; + +template +using VectorChunk = Chunk, 0u, dim_x, dim_y>; + +template +using ChunkStruct = Struct< + Member, + Member, "dfs">, + Member, "dfs_old"> +>; + +template +using VelChunk = Chunk, 0u, dim_x, dim_y>; + +template +using RhoChunk = Chunk, 0u, dim_x, dim_y>; + +template +using MacroStruct = Struct< + Member, "velocity">, + Member, "density"> +>; + +//template +//using ParticleArray = Array< +// Particle +//>; +} + +template void init(){ } +template saw::error_or kmain(int argc, char** argv){ - using namespace kel; + using dfi = df_info; + + (void) argc; + (void) argv; return saw::make_void(); } +} +} + int main(int argc, char** argv){ - auto eov = kmain(argc,argv); + using namespace kel::lbm; + auto eov = kmain(argc,argv); if(eov.is_error()){ auto& err = eov.get_error(); -- cgit v1.2.3 From 5176bfc479a15a5840c4762a9a6c51d5591e41f5 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 13 Jul 2026 15:05:18 +0200 Subject: Prepare for separation --- modules/core/c++/descriptor/d2q9.hpp | 7 +++++++ 1 file changed, 7 insertions(+) create mode 100644 modules/core/c++/descriptor/d2q9.hpp diff --git a/modules/core/c++/descriptor/d2q9.hpp b/modules/core/c++/descriptor/d2q9.hpp new file mode 100644 index 0000000..f675a99 --- /dev/null +++ b/modules/core/c++/descriptor/d2q9.hpp @@ -0,0 +1,7 @@ +#pragma once + +namespace kel { +namespace lbm { + +} +} -- cgit v1.2.3 From 8d115d373795d173acfb98d3115e38c527d7e730 Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 13 Jul 2026 16:46:49 +0200 Subject: Prepare separation into files --- modules/core/c++/component/bgk.hpp | 1 + modules/core/c++/component/bounce_back.hpp | 1 + modules/core/c++/component/no_bounce_back.hpp | 1 + modules/core/c++/component/psm.hpp | 83 ++++++++++++++++++++++++++ modules/core/c++/component/zou_he_density.hpp | 1 + modules/core/c++/component/zou_he_momentum.hpp | 1 + 6 files changed, 88 insertions(+) create mode 100644 modules/core/c++/component/bgk.hpp create mode 100644 modules/core/c++/component/bounce_back.hpp create mode 100644 modules/core/c++/component/no_bounce_back.hpp create mode 100644 modules/core/c++/component/psm.hpp create mode 100644 modules/core/c++/component/zou_he_density.hpp create mode 100644 modules/core/c++/component/zou_he_momentum.hpp diff --git a/modules/core/c++/component/bgk.hpp b/modules/core/c++/component/bgk.hpp new file mode 100644 index 0000000..6f70f09 --- /dev/null +++ b/modules/core/c++/component/bgk.hpp @@ -0,0 +1 @@ +#pragma once diff --git a/modules/core/c++/component/bounce_back.hpp b/modules/core/c++/component/bounce_back.hpp new file mode 100644 index 0000000..6f70f09 --- /dev/null +++ b/modules/core/c++/component/bounce_back.hpp @@ -0,0 +1 @@ +#pragma once diff --git a/modules/core/c++/component/no_bounce_back.hpp b/modules/core/c++/component/no_bounce_back.hpp new file mode 100644 index 0000000..6f70f09 --- /dev/null +++ b/modules/core/c++/component/no_bounce_back.hpp @@ -0,0 +1 @@ +#pragma once diff --git a/modules/core/c++/component/psm.hpp b/modules/core/c++/component/psm.hpp new file mode 100644 index 0000000..cada53c --- /dev/null +++ b/modules/core/c++/component/psm.hpp @@ -0,0 +1,83 @@ +#pragma once + +#include "../macroscopic.hpp" +#include "../equilibrium.hpp" +#include "common.hpp" + +namespace kel { +namespace lbm { +namespace cmpt { +struct PSM {}; +} + +/** + * PSM collision operator for LBM + */ +template +class component { +private: + saw::data relaxation_; + saw::data frequency_; +public: + component( + typename saw::native_data_type::type relaxation__ + ): + relaxation_{relaxation__} + { + saw::data one; + one = 1.0; + frequency_ = one / relaxation_; + } + + template + void apply(const saw::data& field, const saw::data& macros, saw::data> index, saw::data time_step) const { + + using dfi = df_info; + bool is_even = ((time_step.get() % 2) == 0); + + auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">(); + auto& porous_f = macros.template get<"porosity">(); + + auto& rho_f = macros.template get<"density">(); + auto& vel_f = macros.template get<"velocity">(); + + saw::data>& rho = rho_f.at(index); + saw::data>& vel = vel_f.at(index); + + compute_rho_u(dfs_old_f.at(index),rho,vel); + + auto eq = equilibrium(rho,vel); + + saw::data one{1.0}; + auto& porous = porous_f.at(index); + auto flip_porous = one - porous.at({}); + + auto& dfs = dfs_old_f.at(index); + + auto dfs_cpy = dfs; + + for(uint64_t i = 0u; i < Descriptor::Q; ++i){ + uint64_t i_opp = dfi::opposite_index[i]; + dfs.at({i}) = dfs_cpy.at({i}) + frequency_ * (eq.at(i) - dfs_cpy.at({i})) * porous.at({}) + (dfs_cpy.at({i_opp}) - dfs_cpy.at({i}) ) * flip_porous; + } + + auto& force_f = macros.template get<"force">(); + auto& force = force_f.at(index); + for(uint64_t k{0u}; k < Descriptor::D; ++k){ + force.at({{k}}).set(0); + } + + for(uint64_t i{0u}; i < Descriptor::Q; ++i){ + uint64_t i_opp = dfi::opposite_index[i]; + auto dfs_diff = dfs.at({i}) - dfs.at({i_opp}); + for(uint64_t k{0u}; k < Descriptor::D; ++k){ + force.at({{k}}) = force.at({{k}}) + dfs_diff * dfi::directions[i][k]; + } + } + + force = force * porous; + } +}; + +} +} diff --git a/modules/core/c++/component/zou_he_density.hpp b/modules/core/c++/component/zou_he_density.hpp new file mode 100644 index 0000000..6f70f09 --- /dev/null +++ b/modules/core/c++/component/zou_he_density.hpp @@ -0,0 +1 @@ +#pragma once diff --git a/modules/core/c++/component/zou_he_momentum.hpp b/modules/core/c++/component/zou_he_momentum.hpp new file mode 100644 index 0000000..6f70f09 --- /dev/null +++ b/modules/core/c++/component/zou_he_momentum.hpp @@ -0,0 +1 @@ +#pragma once -- cgit v1.2.3 From 0dcc910d33832e2c5c29a6317467f85accf77aac Mon Sep 17 00:00:00 2001 From: "Claudius \"keldu\" Holeksa" Date: Mon, 13 Jul 2026 20:28:07 +0200 Subject: Prep for two way --- examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp b/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp index cae8a4d..1693c5e 100644 --- a/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp +++ b/examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp @@ -180,10 +180,13 @@ saw::error_or step( component> bb; component,encode::Sycl> abb; - component,encode::Sycl> flow_in{1.0}; + component,encode::Sycl> flow_in{1.0001}; component,encode::Sycl> flow_out{1.0}; - component> opa{{},{},eps}; + component> optw; + h.parallel_for(acpp::sycl::range<1u>{1u},[=](acpp::sycl::id<1u> idx){ + + }).wait(); h.parallel_for(acpp::sycl::range{dim_x,dim_y}, [=](acpp::sycl::id idx){ saw::data> index; @@ -294,7 +297,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; } -- cgit v1.2.3