diff options
| -rw-r--r-- | default.nix | 2 | ||||
| -rw-r--r-- | examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp | 2 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/.nix/derivation.nix | 5 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 98 | ||||
| -rw-r--r-- | lib/core/c++/chunk.hpp | 6 | ||||
| -rw-r--r-- | lib/core/c++/geometry.hpp | 4 | ||||
| -rw-r--r-- | lib/core/c++/lbm.hpp | 1 | ||||
| -rw-r--r-- | lib/sycl/c++/SConscript | 9 | ||||
| -rw-r--r-- | lib/sycl/c++/data.hpp | 25 | ||||
| -rw-r--r-- | lib/sycl/c++/lbm.hpp | 3 |
10 files changed, 117 insertions, 38 deletions
diff --git a/default.nix b/default.nix index 461335c..a8ec430 100644 --- a/default.nix +++ b/default.nix @@ -39,7 +39,7 @@ let forstio = (import ((builtins.fetchTarball { url = "https://git.keldu.de/forstio-forstio/snapshot/master.tar.gz"; - sha256 = "sha256:0ckfcygk4f9x51y45hs5jpi95fjdfqfhl6bx6y374f7x98ng9xx3"; + sha256 = "sha256:09xxs9dlnd52k1082lb5p034ik4jfg9zd5zc1x5da2clqhdb6s4m"; }) + "/default.nix"){ inherit stdenv; inherit clang-tools; diff --git a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp index 0f0d219..1aaa5c0 100644 --- a/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp +++ b/examples/poiseulle_2d_gpu/poiseulle_2d_gpu.cpp @@ -338,9 +338,7 @@ saw::error_or<void> kel_main(int argc, char** argv){ } auto& lbm_dir = eo_lbm_dir.get_value(); auto out_dir = lbm_dir / "poiseulle_channel_2d_gpu"; - std::filesystem::create_directories(out_dir); - // lbm::converter<lbm::sch::Float64> conv { // delta_x diff --git a/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix b/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix index 127243d..e08f42a 100644 --- a/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix +++ b/examples/poiseulle_particles_2d_gpu/.nix/derivation.nix @@ -26,11 +26,12 @@ stdenv.mkDerivation { forstio.async forstio.codec forstio.codec-unit - forstio.remote + forstio.remote + forstio.remote-filesystem forstio.codec-json adaptive-cpp kel.lbm.core -# kel-lbm.sycl + kel.lbm.sycl ]; preferLocalBuild = true; diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index 8bc60c9..be6b4f0 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -1,8 +1,15 @@ #include <kel/lbm/lbm.hpp> -#include <AdaptiveCpp/sycl/sycl.hpp> +#include <kel/lbm/sycl/lbm.hpp> + +#include <forstio/remote/filesystem/easy.hpp> +#include <forstio/codec/json/json.hpp> namespace kel { namespace lbm { + +constexpr uint64_t dim_x = 16u; +constexpr uint64_t dim_y = 4u; + namespace sch { using namespace saw::schema; @@ -14,6 +21,63 @@ using CellStruct = Struct< Member<Vector<T,Desc::D>, "velocity">, Member<Vector<T,Desc::D>, "force"> >; + + +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 ChunkStruct = Struct< + Member<InfoChunk, "info"> +>; + +} + +template<typename T, typename Desc> +saw::error_or<void> setup_initial_conditions(saw::data<sch::ChunkStruct<T,Desc>>& fields){ + auto& info_f = fields.get<"info">(); + // Set everything as walls + iterator<Desc::D>::apply( + [&](auto& index){ + info_f.at(index).set(2u); + }, + {{0u,0u}}, + {{dim_x, dim_y}}, + {{0u,0u}} + ); + // Fluid + iterator<Desc::D>::apply( + [&](auto& index){ + info_f.at(index).set(1u); + }, + {{0u,0u}}, + {{dim_x, dim_y}}, + {{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}} + ); + + return saw::make_void(); } template<typename T, typename Desc> @@ -25,48 +89,54 @@ saw::error_or<void> step(){ } template<typename T, typename Desc> -saw::error_or<void> kel_main(int argc, char** argv){ - using namespace kel; +saw::error_or<void> lbm_main(int argc, char** argv){ + using namespace kel::lbm; - using dfi = lbm::df_info<T,Desc>; + using dfi = df_info<T,Desc>; - auto eo_lbm_dir = lbm::output_directory(); + 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"; - lbm::converter<lbm::sch::Float64> conv { + converter<sch::Float64> conv { // delta_x {{1.0}}, // delta_t {{1.0}} }; - uint64_t x_d = 256u; - uint64_t y_d = 64u; - saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{x_d,y_d}}; - saw::data<sch::Array<lbm::sch::CellStruct<T,Desc>,Desc::D>> lbm_data{meta}; + // saw::data<sch::FixedArray<sch::UInt64,Desc::D>> meta{{dim_x,dim_y}}; + saw::data<sch::ChunkStruct<T,Desc>> lbm_data{}; acpp::sycl::queue sycl_q; sycl_q.wait(); { - auto eov = setup_initial_conditions(lbm_data); + auto eov = setup_initial_conditions<T,Desc>(lbm_data); if(eov.is_error()){ return eov; } } - + /* + iterator<Desc::D>::apply( + [&](auto& index){ + std::cout<<index.at({0u}).get()<<" "<<index.at({1u}).get()<<" "<<lbm_data.get<"info">().at(index).template cast_to<sch::UInt16>().get()<<"\n"; + }, + {{0u,0u}}, + {{dim_x, dim_y}} + ); + */ return saw::make_void(); } -using FloatT = kel::sch::Float32; +using FloatT = kel::lbm::sch::Float32; int main(int argc, char** argv){ - auto eov = kel_main<FloatT,kel::lbm::sch::D2Q9>(argc, 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(); diff --git a/lib/core/c++/chunk.hpp b/lib/core/c++/chunk.hpp index 607bec0..ad7de34 100644 --- a/lib/core/c++/chunk.hpp +++ b/lib/core/c++/chunk.hpp @@ -76,4 +76,10 @@ public: return data<schema::FixedArray<schema::UInt64, sizeof...(Sides)>,Encode>{{Sides...}}; } }; + +template<typename Sch, uint64_t Ghost, uint64_t... Sides> +struct meta_schema<kel::lbm::sch::Chunk<Sch,Ghost,Sides...>> { + using MetaSchema = typename meta_schema<Sch>::MetaSchema; +}; + } diff --git a/lib/core/c++/geometry.hpp b/lib/core/c++/geometry.hpp index 6875e90..c8a48a6 100644 --- a/lib/core/c++/geometry.hpp +++ b/lib/core/c++/geometry.hpp @@ -21,7 +21,9 @@ class component<Schema, Desc, cmpt::PoiseulleChannel, Encode> final { private: public: template<typename CellFieldSchema> - void apply(saw::data<CellFieldSchema,Encode>& field, const saw::data<sch::FixedArraysch::UInt64,Desc::D>) + void apply(saw::data<CellFieldSchema,Encode>& field, const saw::data<sch::FixedArraysch::UInt64,Desc::D>){ + auto& info_f = field.template get<"info">(); + } }; // Ghost - 0 diff --git a/lib/core/c++/lbm.hpp b/lib/core/c++/lbm.hpp index aff38e9..00f153a 100644 --- a/lib/core/c++/lbm.hpp +++ b/lib/core/c++/lbm.hpp @@ -2,6 +2,7 @@ #include "schema.hpp" #include "flatten.hpp" +#include "chunk.hpp" #include "descriptor.hpp" #include "boundary.hpp" #include "converter.hpp" diff --git a/lib/sycl/c++/SConscript b/lib/sycl/c++/SConscript index 85a078f..2ed63ba 100644 --- a/lib/sycl/c++/SConscript +++ b/lib/sycl/c++/SConscript @@ -15,18 +15,13 @@ core_env = env.Clone(); core_env.sources = sorted(glob.glob(dir_path + "/*.cpp")); core_env.headers = sorted(glob.glob(dir_path + "/*.hpp")); -core_env.particle_headers = sorted(glob.glob(dir_path + "/particle/*.hpp")); -core_env.particle_geometry_headers = sorted(glob.glob(dir_path + "/particle/geometry/*.hpp")); - env.sources += core_env.sources; env.headers += core_env.headers; ## Static lib objects = [] core_env.add_source_files(objects, core_env.sources, shared=False); -env.library_static = core_env.StaticLibrary('#build/kel-lbm', [objects]); +env.library_static = core_env.StaticLibrary('#build/kel-lbm-sycl', [objects]); env.Install('$prefix/lib/', env.library_static); -env.Install('$prefix/include/kel/lbm/', core_env.headers); -env.Install('$prefix/include/kel/lbm/particle/', core_env.particle_headers); -env.Install('$prefix/include/kel/lbm/particle/geometry/', core_env.particle_geometry_headers); +env.Install('$prefix/include/kel/lbm/sycl/', core_env.headers); diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index 44bc5dc..bb8b4bf 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -1,6 +1,7 @@ #pragma once #include "common.hpp" +#include <kel/lbm/lbm.hpp> namespace kel { namespace lbm { @@ -19,9 +20,9 @@ struct struct_has_only_equal_dimension_array{}; namespace saw { template<uint64_t Ghost, uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode> -class data<schema::Struct<schema::Member<schema::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode>> final { +class data<schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode> > final { public: - static constexpr data<schema::FixedArray meta = {{Meta...}}; + static constexpr data<schema::FixedArray<schema::UInt64,sizeof...(Meta)>> meta = {{Meta...}}; using StorageT = std::tuple<data<Sch,Encode>*...>; private: @@ -50,14 +51,14 @@ public: if(not arg){ return; } - sycl::free(arg,*q_); + acpp::sycl::free(arg,*q_); arg = nullptr; },members_); } template<saw::string_literal K> auto* get_ptr(){ - return std::get<parameter_key_pack_index<K, Keys...>::value(members_); + return std::get<parameter_key_pack_index<K, Keys...>::value>(members_); } template<saw::string_literal K> @@ -67,6 +68,7 @@ public: return *ptr; } }; +} namespace kel { namespace lbm { @@ -76,14 +78,15 @@ struct sycl_malloc_struct_helper; template<typename... Members, typename Encode> struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final { + using Schema = sch::Struct<Members...>; + template<uint64_t i> - static saw::error_or<void> allocate_on_device_member(typename data<Sch,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){ + static saw::error_or<void> allocate_on_device_member(typename saw::data<typename saw::parameter_pack_type<i,Members...>::type::ValueType,encode::Sycl<Encode>>::StorageT& storage, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ using M = typename saw::parameter_pack_type<i,Members...>::type; auto& ptr = std::get<i>(storage); - - ptr = sycl::malloc_device<M::ValueType::InnerSchema>(,q); + ptr = sycl::malloc_device<M::ValueType::InnerSchema>(1u,q); return allocate_on_device_member<i+1u>(storage,q); } @@ -91,8 +94,8 @@ struct sycl_malloc_struct_helper<sch::Struct<Members...>, Encode> final { return saw::make_void(); } - static saw::error_or<void> allocate_on_device(data<Sch,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ - typename data<Sch,encode::Sycl<Encode>>::StorageT storage; + static saw::error_or<void> allocate_on_device(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ + typename saw::data<Schema,encode::Sycl<Encode>>::StorageT storage; auto eov = allocate_on_device_member<0u>(storage,q); sycl_data = {storage, q}; @@ -112,8 +115,8 @@ public: ~device() = default; template<typename Sch, typename Encode> - saw::error_or<void> allocate_on_device(data<Sch,encode::Sycl<Encode>>& sycl_data){ - auto eov = sycl_malloc_struct_helper<Sch,Encode>::allocate_on_device(sycl_data, q_); + saw::error_or<void> allocate_on_device(saw::data<Sch,encode::Sycl<Encode>>& sycl_data){ + auto eov = impl::sycl_malloc_struct_helper<Sch,Encode>::allocate_on_device(sycl_data, q_); if(eov.is_error()){ return eov; } diff --git a/lib/sycl/c++/lbm.hpp b/lib/sycl/c++/lbm.hpp new file mode 100644 index 0000000..3e67ae6 --- /dev/null +++ b/lib/sycl/c++/lbm.hpp @@ -0,0 +1,3 @@ +#pragma once + +#include "data.hpp" |