summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
Diffstat
-rw-r--r--examples/moving_poiseulle_particles_2d_hlbm_gpu/sim.cpp22
-rw-r--r--examples/poiseulle_moving_particle_2d_psm_gpu/sim.cpp9
-rw-r--r--modules/core/c++/component/bgk.hpp1
-rw-r--r--modules/core/c++/component/bounce_back.hpp1
-rw-r--r--modules/core/c++/component/no_bounce_back.hpp1
-rw-r--r--modules/core/c++/component/psm.hpp83
-rw-r--r--modules/core/c++/component/zou_he_density.hpp1
-rw-r--r--modules/core/c++/component/zou_he_momentum.hpp1
-rw-r--r--modules/core/c++/descriptor/d2q9.hpp7
-rw-r--r--poc/mixed_precision/sim.cpp24
-rw-r--r--poc/triangle_stream/.nix/derivation.nix41
-rw-r--r--poc/triangle_stream/SConscript34
-rw-r--r--poc/triangle_stream/SConstruct81
-rw-r--r--poc/triangle_stream/sim.cpp80
14 files changed, 361 insertions, 25 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<void> setup_initial_conditions(
df_f.get_dims(),
{{1u,1u}}
);
- /* Technically the particle group should handle this case
- 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()
- );
- */
-
-
return saw::make_void();
}
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<void> step(
component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
component<T,Desc,cmpt::AntiBounceBack<0u>,encode::Sycl<saw::encode::Native>> abb;
- component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{1.0};
+ component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{1.0001};
component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
- component<T,Desc,cmpt::OneParticleAt, encode::Sycl<saw::encode::Native>> opa{{},{},eps};
+ component<T,Desc,cmpt::OneParticleTwoWay, encode::Sycl<saw::encode::Native>> optw;
+ h.parallel_for(acpp::sycl::range<1u>{1u},[=](acpp::sycl::id<1u> idx){
+
+ }).wait();
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;
@@ -294,7 +297,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
sycl_q.wait();
{
- auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr);
+ auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr);
if(eov.is_error()){
return eov;
}
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<typename T, typename Descriptor, typename Encode>
+class component<T, Descriptor, cmpt::PSM, Encode> {
+private:
+ saw::data<T> relaxation_;
+ saw::data<T> frequency_;
+public:
+ component(
+ typename saw::native_data_type<T>::type relaxation__
+ ):
+ relaxation_{relaxation__}
+ {
+ saw::data<T> one;
+ one = 1.0;
+ frequency_ = one / relaxation_;
+ }
+
+ template<typename CellFieldSchema, typename MacroFieldSchema>
+ void apply(const saw::data<CellFieldSchema, Encode>& field, const saw::data<MacroFieldSchema,Encode>& macros, saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> index, saw::data<sch::UInt64> time_step) const {
+
+ using dfi = df_info<T,Descriptor>;
+ 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<sch::Scalar<T>>& rho = rho_f.at(index);
+ saw::data<sch::Vector<T,Descriptor::D>>& vel = vel_f.at(index);
+
+ compute_rho_u<T,Descriptor>(dfs_old_f.at(index),rho,vel);
+
+ auto eq = equilibrium<T,Descriptor>(rho,vel);
+
+ saw::data<T> 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
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 {
+
+}
+}
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<void> 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]: "<<err.get_category();
+ auto err_msg = err.get_message();
+ if(not err_msg.empty()){
+ std::cerr<<" - "<<err_msg;
+ }
+ std::cerr<<std::endl;
+ return err.code();
+ }
+ return 0;
+}
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
new file mode 100644
index 0000000..36986e4
--- /dev/null
+++ b/poc/triangle_stream/sim.cpp
@@ -0,0 +1,80 @@
+#include <kel/lbm/lbm.hpp>
+
+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<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">
+>;
+
+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">
+>;
+
+//template<typename T, typename Desc>
+//using ParticleArray = Array<
+// Particle<T,Desc::D>
+//>;
+}
+
+template<typename T, typename Desc>
+void init(){
+}
+
+template<typename T, typename Desc>
+saw::error_or<void> kmain(int argc, char** argv){
+ using dfi = df_info<T,Desc>;
+
+ (void) argc;
+ (void) argv;
+
+ return saw::make_void();
+}
+
+}
+}
+
+int main(int argc, char** argv){
+ using namespace kel::lbm;
+ auto eov = kmain<sch::Float32,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(not err_msg.empty()){
+ std::cerr<<" - "<<err_msg;
+ }
+ std::cerr<<std::endl;
+ return err.code();
+ }
+ return 0;
+}
generated by cgit v1.2.3 (git 2.25.1) at 2026-07-21 15:35:07 +0000