diff options
| author | Claudius "keldu" Holeksa <mail@keldu.de> | 2026-03-06 21:05:36 +0100 |
|---|---|---|
| committer | Claudius "keldu" Holeksa <mail@keldu.de> | 2026-03-06 21:05:36 +0100 |
| commit | 25c85bf962e0646f8e03f67fd4982450f41ee6a6 (patch) | |
| tree | 1310dca3d2a9e372ead13d166a57db8cfcbbfe84 /examples/poiseulle_particles_2d_gpu | |
| parent | 507440cdf786f7a1a83bdf22371f1844910382fb (diff) | |
| download | libs-lbm-dev.tar.gz | |
Work finished for this weekdev
Diffstat (limited to 'examples/poiseulle_particles_2d_gpu')
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/.nix/derivation.nix | 41 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/SConscript | 34 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/SConstruct | 81 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 469 |
4 files changed, 0 insertions, 625 deletions
diff --git a/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix b/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix deleted file mode 100644 index 517b6b8..0000000 --- a/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix +++ /dev/null @@ -1,41 +0,0 @@ -{ lib -, stdenv -, scons -, clang-tools -, forstio -, python3 -, pname -, version -, adaptive-cpp -, kel -}: - -stdenv.mkDerivation { - pname = pname + "-examples-" + "poiseulle_2d_gpu"; - 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/examples/poiseulle_particles_2d_gpu/SConscript b/examples/poiseulle_particles_2d_gpu/SConscript deleted file mode 100644 index 5cad56f..0000000 --- a/examples/poiseulle_particles_2d_gpu/SConscript +++ /dev/null @@ -1,34 +0,0 @@ -#!/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/poiseulle_particles_2d_gpu', [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/examples/poiseulle_particles_2d_gpu/SConstruct b/examples/poiseulle_particles_2d_gpu/SConstruct deleted file mode 100644 index 0611b67..0000000 --- a/examples/poiseulle_particles_2d_gpu/SConstruct +++ /dev/null @@ -1,81 +0,0 @@ -#!/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/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp deleted file mode 100644 index 644e4d1..0000000 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ /dev/null @@ -1,469 +0,0 @@ -#include <kel/lbm/sycl/lbm.hpp> -#include <kel/lbm/lbm.hpp> -#include <kel/lbm/particle.hpp> - -#include <forstio/io/io.hpp> -#include <forstio/remote/filesystem/easy.hpp> -#include <forstio/codec/json/json.hpp> -#include <forstio/codec/simple.hpp> - -namespace kel { -namespace lbm { - -constexpr uint64_t dim_y = 256ul; -constexpr uint64_t dim_x = dim_y * 20ul; - -constexpr uint64_t particle_amount = 1ul; - -namespace sch { -using namespace saw::schema; - -using InfoChunk = Chunk<UInt8, 0u, dim_x, dim_y>; - -template<typename T, typename Desc> -using DfChunk = Chunk<FixedArray<T,Desc::Q>, 1u, dim_x, dim_y>; - -template<typename T, typename Desc> -using ScalarChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>; - -template<typename T, typename Desc> -using VectorChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>; - -template<typename T, typename Desc> -using ChunkStruct = Struct< - Member<InfoChunk, "info">, - Member<DfChunk<T,Desc>, "dfs">, - Member<DfChunk<T,Desc>, "dfs_old">, - Member<VectorChunk<T,Desc>, "particle_N">, - Member<ScalarChunk<T,Desc>, "particle_D"> ->; - -template<typename T, typename Desc> -using VelChunk = Chunk<Vector<T,Desc::D>, 0u, dim_x, dim_y>; - -template<typename T> -using RhoChunk = Chunk<Scalar<T>, 0u, dim_x, dim_y>; - -template<typename T, typename Desc> -using MacroStruct = Struct< - Member<VelChunk<T,Desc>, "velocity">, - Member<RhoChunk<T>, "density">, - Member<ScalarChunk<T,Desc>, "porosity"> ->; - -//template<typename T, typename Desc> -//using ParticleArray = Array< -// Particle<T,Desc::D> -//>; -} - -template<typename T, typename Desc> -saw::error_or<void> setup_initial_conditions( - saw::data<sch::ChunkStruct<T,Desc>>& fields, - saw::data<sch::MacroStruct<T,Desc>>& macros, - saw::data<sch::FixedArray<sch::Particle<T,Desc::D>, particle_amount>>& particles -){ - auto& info_f = fields.template get<"info">(); - auto& porous_f = macros.template get<"porosity">(); - // Set everything as walls - iterator<Desc::D>::apply( - [&](auto& index){ - info_f.at(index).set(1u); - }, - {}, - info_f.get_dims(), - {} - ); - // Fluid - iterator<Desc::D>::apply( - [&](auto& index){ - info_f.at(index).set(2u); - }, - {}, - info_f.get_dims(), - {{1u,1u}} - ); - - // Inflow - iterator<Desc::D>::apply( - [&](auto& index){ - info_f.at(index).set(3u); - }, - {{0u,0u}}, - {{1u,dim_y}}, - {{0u,1u}} - ); - - // Outflow - iterator<Desc::D>::apply( - [&](auto& index){ - info_f.at(index).set(4u); - }, - {{dim_x-1u,0u}}, - {{dim_x, dim_y}}, - {{0u,1u}} - ); - // - auto& df_f = fields.template get<"dfs_old">(); - auto& rho_f = macros.template get<"density">(); - auto& vel_f = macros.template get<"velocity">(); - auto& por_f = macros.template get<"porosity">(); - - iterator<Desc::D>::apply( - [&](auto& index){ - auto& df = df_f.at(index); - auto& rho = rho_f.at(index); - por_f.at(index).at({}) = {1}; - rho.at({}) = {1}; - auto& vel = vel_f.at(index); - auto eq = equilibrium<T,Desc>(rho,vel); - - df = eq; - }, - {},// 0-index - df_f.get_dims() - ); - - iterator<Desc::D>::apply( - [&](auto& index){ - auto& df = df_f.at(index); - auto& rho = rho_f.at(index); - rho.at({}) = {1}; - auto& vel = vel_f.at(index); - if(info_f.at(index).get() == 2u){ - vel.at({{0u}}) = 0.0; - } - auto eq = equilibrium<T,Desc>(rho,vel); - - df = eq; - }, - {},// 0-index - df_f.get_dims(), - {{1u,1u}} - ); - - iterator<Desc::D>::apply( - [&](auto& index){ - saw::data<sch::Vector<T,Desc::D>> 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<T>(); - ind_vec.at({{1u}}) = index.at({{1u}}).template cast_to<T>(); - - 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() - ); - - for(saw::data<sch::UInt64> i{0u}; i < saw::data<sch::UInt64>{particle_amount}; ++i){ - auto& part = particles.at(i); - - saw::data<sch::Vector<T,Desc::D>> pos; - pos.at({{0u}}) = dim_x * 0.25; - pos.at({{1u}}) = dim_y * 0.5; - saw::data<sch::Scalar<T>> rad, dense, dt; - rad.at({}) = dim_y * 0.1; - dense.at({}) = 1.0; - dt.at({}) = 1.0; - part = create_spheroid_particle( - pos,{},{}, - {},{},{}, - rad, dense,dt - ); - } - - return saw::make_void(); -} - -template<typename T, typename Desc> -saw::error_or<void> step( - saw::data<sch::Ptr<sch::ChunkStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& fields, - saw::data<sch::Ptr<sch::MacroStruct<T,Desc>>,encode::Sycl<saw::encode::Native>>& macros, - saw::data<sch::FixedArray<sch::Particle<T,Desc::D>, particle_amount>>& particles, - saw::data<sch::UInt64> t_i, - device& dev -){ - auto& q = dev.get_handle(); - auto& info_f = fields.template get<"info">(); - auto& porous_f = macros.template get<"porosity">(); - - { - auto& p = particles.at({{0u}}); - - auto& p_coll = p.template get<"collision">(); - auto& p_rad = p_coll.template get<"radius">(); - } - - - // auto coll_ev = - q.submit([&](acpp::sycl::handler& h){ - // Need nicer things to handle the flow. I see improvement here - component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{0.65}; - // component<T,Desc,cmpt::HLBM,encode::Sycl<saw::encode::Native>> collision{0.65}; - // component<T,Desc,cmpt::PSM,encode::Sycl<saw::encode::Native>> collision{0.65}; - component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb; - - saw::data<sch::Scalar<T>> rho_b; - rho_b.at({}) = 1.0; - saw::data<sch::Vector<T,Desc::D>> vel_b; - vel_b.at({{0u}}) = 0.015; - - component<T,Desc,cmpt::Equilibrium,encode::Sycl<saw::encode::Native>> equi{rho_b,vel_b}; - - component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{ - [&](){ - uint64_t target_t_i = 64u; - if(t_i.get() < target_t_i){ - return 1.0 + (0.0015 / target_t_i) * t_i.get(); - } - return 1.0015; - }() - }; - component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0}; - - - h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){ - saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index; - for(uint64_t i = 0u; i < Desc::D; ++i){ - index.at({{i}}).set(idx[i]); - } - - auto info = info_f.at(index); - - switch(info.get()){ - case 0u: - break; - case 1u: - bb.apply(fields,index,t_i); - break; - case 2u: - collision.apply(fields,macros,index,t_i); - break; - case 3u: - flow_in.apply(fields,index,t_i); - collision.apply(fields,macros,index,t_i); - break; - case 4u: - flow_out.apply(fields,index,t_i); - collision.apply(fields,macros,index,t_i); - break; - default: - break; - } - }); - }).wait(); - - - // Step - /* - q.submit([&](acpp::sycl::handler& h){ - // h.depends_on(collision_ev); - }).wait(); - */ - - return saw::make_void(); -} -} -} - -template<typename T, typename Desc> -saw::error_or<void> lbm_main(int argc, char** argv){ - using namespace kel::lbm; - - using dfi = df_info<T,Desc>; - - auto eo_lbm_dir = output_directory(); - if(eo_lbm_dir.is_error()){ - 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<saw::err::critical>("Could not create output directory"); - } - } - - converter<T> conv { - // delta_x - {{1.0}}, - // delta_t - {{1.0}} - }; - - print_lbm_meta<T,Desc>(conv,{0.05},{0.01},{0.4 * dim_y}); - - // saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}}; - auto lbm_data_ptr = saw::heap<saw::data<sch::ChunkStruct<T,Desc>>>(); - auto lbm_macro_data_ptr = saw::heap<saw::data<sch::MacroStruct<T,Desc>>>(); - auto lbm_particle_data_ptr = saw::heap<saw::data<sch::FixedArray<sch::Particle<T,Desc::D>, particle_amount>>>(); - - std::cout<<"Estimated Bytes: "<<memory_estimate<sch::ChunkStruct<T,Desc>,sch::MacroStruct<T,Desc>>().get()<<std::endl; - - auto eo_aio = saw::setup_async_io(); - if(eo_aio.is_error()){ - return std::move(eo_aio.get_error()); - } - auto& aio = eo_aio.get_value(); - saw::wait_scope wait{aio.event_loop}; - - bool krun = true; - bool print_status = false; - aio.event_port.on_signal(saw::Signal::Terminate).then([&](){ - krun = false; - }).detach(); - aio.event_port.on_signal(saw::Signal::User1).then([&](){ - print_status = true; - }).detach(); - - device dev; - - auto& sycl_q = dev.get_handle(); - - sycl_q.wait(); - { - auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr); - if(eov.is_error()){ - return eov; - } - } - { - auto eov = write_vtk_file(out_dir,"initial_state",0u,*lbm_data_ptr); - if(eov.is_error()){ - return eov; - } - } - - saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q}; - saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_data{sycl_q}; - saw::data<sch::FixedArray<sch::Particle<T,Desc::D>, particle_amount>, encode::Sycl<saw::encode::Native>> 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); - { - auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data); - if(eov.is_error()){ - return eov; - } - } - { - auto eov = dev.copy_to_device(*lbm_macro_data_ptr,lbm_sycl_macro_data); - if(eov.is_error()){ - 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<sch::UInt64> time_steps{4096ul}; - - auto& info_f = lsd_view.template get<"info">(); - - for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){ - // BC + Collision - { - auto eov = step<T,Desc>(lsd_view,lsdm_view,*lbm_particle_data_ptr,i,dev); - if(eov.is_error()){ - return eov; - } - } - sycl_q.wait(); - { - { - auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - { - auto eov = write_vtk_file(out_dir,"m",i.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - } - /* - { - { - auto eov = dev.copy_to_host(lbm_sycl_data,*lbm_data_ptr); - if(eov.is_error()){ - return eov; - } - } - { - auto eov = write_csv_file(out_dir,"lbm",i.get(), *lbm_data_ptr); - if(eov.is_error()){ - return eov; - } - } - } - */ - // Stream - sycl_q.submit([&](acpp::sycl::handler& h){ - component<T,Desc,cmpt::Stream,encode::Sycl<saw::encode::Native>> stream; - - h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){ - saw::data<sch::FixedArray<sch::UInt64,Desc::D>> index; - for(uint64_t i = 0u; i < Desc::D; ++i){ - index.at({{i}}).set(idx[i]); - } - - auto info = info_f.at(index); - - if(info.get() > 0u){ - stream.apply(lsd_view,index,i); - } - }); - }).wait(); - wait.poll(); - if(print_status){ - std::cout<<"Status: "<<i.get()<<" of "<<time_steps.get()<<" - "<<(i.template cast_to<sch::Float64>().get() * 100 / time_steps.get())<<"%"<<std::endl; - print_status = false; - } - print_progress_bar(i.get(), time_steps.get()-1u); - } - - // After Loop - sycl_q.wait(); - { - auto eov = write_vtk_file(out_dir,"m",time_steps.get(), *lbm_macro_data_ptr); - if(eov.is_error()){ - return eov; - } - } - - sycl_q.wait(); - return saw::make_void(); -} - -using FloatT = kel::lbm::sch::Float32; - -int main(int argc, char** argv){ - auto eov = lbm_main<FloatT,kel::lbm::sch::D2Q9>(argc, argv); - if(eov.is_error()){ - auto& err = eov.get_error(); - std::cerr<<"[Error] "<<err.get_category(); - auto err_msg = err.get_message(); - if(err_msg.size() > 0u){ - std::cerr<<" - "<<err_msg; - } - std::cerr<<std::endl; - return err.get_id(); - } - return 0; -} |
