summaryrefslogtreecommitdiff
path: root/examples
diff options
context:
space:
mode:
Diffstat (limited to 'examples')
-rw-r--r--examples/poiseulle_3d_gpu/sim.cpp4
-rw-r--r--examples/poiseulle_particles_2d_bgk_gpu/sim.cpp23
-rw-r--r--examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp20
-rw-r--r--examples/poiseulle_particles_2d_psm_gpu/sim.cpp19
-rw-r--r--examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp3
-rw-r--r--examples/settling_cubes_2d_ibm_gpu/sim.cpp189
6 files changed, 195 insertions, 63 deletions
diff --git a/examples/poiseulle_3d_gpu/sim.cpp b/examples/poiseulle_3d_gpu/sim.cpp
index 93df1f8..6916ef2 100644
--- a/examples/poiseulle_3d_gpu/sim.cpp
+++ b/examples/poiseulle_3d_gpu/sim.cpp
@@ -315,8 +315,8 @@ saw::error_or<void> lbm_main(int argc, char** argv){
saw::data<sch::FixedArray<sch::Particle<T,Desc::D>, particle_size>, 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 lsd_view = make_view(lbm_sycl_data);
+ auto lsdm_view = make_view(lbm_sycl_macro_data);
{
auto eov = dev.copy_to_device(*lbm_data_ptr,lbm_sycl_data);
if(eov.is_error()){
diff --git a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp
index 7d1a624..b8ac35e 100644
--- a/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_bgk_gpu/sim.cpp
@@ -183,11 +183,11 @@ saw::error_or<void> step(
component<T,Desc,cmpt::ZouHeHorizontal<true>,encode::Sycl<saw::encode::Native>> flow_in{
[&](){
- uint64_t target_t_i = 64u;
+ uint64_t target_t_i = 32u;
if(t_i.get() < target_t_i){
- return 1.0 + (0.0015 / target_t_i) * t_i.get();
+ return 1.0 + (0.0002 / target_t_i) * t_i.get();
}
- return 1.0015;
+ return 1.0002;
}()
};
component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
@@ -205,19 +205,18 @@ saw::error_or<void> step(
case 0u:
break;
case 1u:
- abb.apply(fields,index,t_i);
+ 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);
- equi.apply(fields,index,t_i);
+ flow_in.apply(fields,index,t_i);
+ //equi.apply(fields,index,t_i);
collision.apply(fields,macros,index,t_i);
break;
case 4u:
- // flow_out.apply(fields,index,t_i);
- equi.apply(fields,index,t_i);
+ flow_out.apply(fields,index,t_i);
collision.apply(fields,macros,index,t_i);
break;
default:
@@ -268,7 +267,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
{{1.0}}
};
- print_lbm_meta<T,Desc>(conv,{0.05},{0.01},{0.4 * dim_y});
+ print_lbm_meta<T,Desc>(conv,{0.1},{1e-4},{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>>>();
@@ -314,8 +313,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_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()){
@@ -329,6 +326,8 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
}
sycl_q.wait();
+ auto lsd_view = make_view(lbm_sycl_data);
+ auto lsdm_view = make_view(lbm_sycl_macro_data);
saw::data<sch::UInt64> time_steps{16u*4096ul};
auto& info_f = lsd_view.template get<"info">();
@@ -343,7 +342,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
sycl_q.wait();
/*
- {
+ if(i.get() % 128u == 0u){
{
auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
if(eov.is_error()){
diff --git a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp
index 9eeda3c..7cfd567 100644
--- a/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_hlbm_gpu/sim.cpp
@@ -191,9 +191,9 @@ saw::error_or<void> step(
[&](){
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.0 + (0.0002 / target_t_i) * t_i.get();
}
- return 1.0015;
+ return 1.0002;
}()
};
component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
@@ -211,19 +211,19 @@ saw::error_or<void> step(
case 0u:
break;
case 1u:
- abb.apply(fields,index,t_i);
+ 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);
- equi.apply(fields,index,t_i);
+ flow_in.apply(fields,index,t_i);
+ // equi.apply(fields,index,t_i);
collision.apply(fields,macros,index,t_i);
break;
case 4u:
- // flow_out.apply(fields,index,t_i);
- equi.apply(fields,index,t_i);
+ flow_out.apply(fields,index,t_i);
+ // equi.apply(fields,index,t_i);
collision.apply(fields,macros,index,t_i);
break;
default:
@@ -274,7 +274,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
{{1.0}}
};
- print_lbm_meta<T,Desc>(conv,{0.05},{0.01},{0.4 * dim_y});
+ print_lbm_meta<T,Desc>(conv,{0.1},{1e-4},{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>>>();
@@ -320,8 +320,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_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()){
@@ -335,6 +333,8 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
}
sycl_q.wait();
+ auto lsd_view = make_view(lbm_sycl_data);
+ auto lsdm_view = make_view(lbm_sycl_macro_data);
saw::data<sch::UInt64> time_steps{16u*4096ul};
auto& info_f = lsd_view.template get<"info">();
diff --git a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp
index 7e76da6..75d1a14 100644
--- a/examples/poiseulle_particles_2d_psm_gpu/sim.cpp
+++ b/examples/poiseulle_particles_2d_psm_gpu/sim.cpp
@@ -207,9 +207,9 @@ saw::error_or<void> step(
[&](){
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.0 + (0.0002 / target_t_i) * t_i.get();
}
- return 1.0015;
+ return 1.0002;
}()
};
component<T,Desc,cmpt::ZouHeHorizontal<false>,encode::Sycl<saw::encode::Native>> flow_out{1.0};
@@ -227,7 +227,7 @@ saw::error_or<void> step(
case 0u:
break;
case 1u:
- abb.apply(fields,index,t_i);
+ bb.apply(fields,index,t_i);
break;
case 2u:
collision.apply(fields,macros,index,t_i);
@@ -289,12 +289,12 @@ saw::error_or<void> lbm_main(int argc, char** argv){
converter<T> conv {
// delta_x
- {{1e-3}},
+ {{1.0}},
// delta_t
- {{1e-06}}
+ {{1.0}}
};
- print_lbm_meta<T,Desc>(conv,{0.1},{1e-2},{0.4 * dim_y *1e-3});
+ print_lbm_meta<T,Desc>(conv,{0.1},{1e-4},{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>>>();
@@ -340,8 +340,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
saw::data<sch::MacroStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_macro_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()){
@@ -355,6 +353,9 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
}
sycl_q.wait();
+ auto lsd_view = make_view(lbm_sycl_data);
+ auto lsdm_view = make_view(lbm_sycl_macro_data);
+
saw::data<sch::UInt64> time_steps{16u*4096ul};
auto& info_f = lsd_view.template get<"info">();
@@ -368,6 +369,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
}
sycl_q.wait();
+ /*
if(i.get() % 32u == 0u){
{
auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
@@ -382,6 +384,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
}
}
+ */
/*
{
{
diff --git a/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp b/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp
index 7ac663f..1e6e75f 100644
--- a/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp
+++ b/examples/poiseulle_particles_channel_2d/poiseulle_particles_channel_2d.cpp
@@ -461,6 +461,8 @@ void couple_particles_to_lattice(
p_acc = p_acc + p_pos_rel_vec;
}
+
+
/* So I want to include the relative velocity from both particles, but this is a bit hard considering I just assumed 0 velocity from the other party
else if(n_macro_cell_particle.get() != (i.get() + 1u) and n_macro_cell_particle.get() > 0u){
// Generally compare
@@ -478,6 +480,7 @@ void couple_particles_to_lattice(
auto& n_info = n_cell.template get<"info">()({0u});
auto& n_macro_cell = macros.at(n_p_cell_pos);
auto& n_macro_cell_particle = n_macro_cell.template get<"particle">();
+ l
// If neighbour is wall, then add force pushing the particle away
if(n_info.get() <= 1u or (n_macro_cell_particle.get() != (i.get()+1u) and n_macro_cell_particle.get() > 0u) ) {
diff --git a/examples/settling_cubes_2d_ibm_gpu/sim.cpp b/examples/settling_cubes_2d_ibm_gpu/sim.cpp
index 3281239..d7b402a 100644
--- a/examples/settling_cubes_2d_ibm_gpu/sim.cpp
+++ b/examples/settling_cubes_2d_ibm_gpu/sim.cpp
@@ -6,6 +6,7 @@
#include <forstio/remote/filesystem/easy.hpp>
#include <forstio/codec/json/json.hpp>
#include <forstio/codec/simple.hpp>
+#include <forstio/codec/math.hpp>
namespace kel {
namespace lbm {
@@ -50,18 +51,25 @@ using MacroStruct = Struct<
>;
template<typename T, typename Desc>
+using ParticleSpheroidGroup = ParticleGroup<
+ T,
+ Desc::D,
+ sch::ParticleCollisionSpheroid<T,2.0f>
+>;
+
+template<typename T, typename Desc>
using ParticleGroups = Tuple<
- ParticleGroup<
- T,Desc::D,sch::ParticleCollisionSpheroid<T>
- >
+ ParticleSpheroidGroup<T,Desc>
>;
+
+
}
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::ParticleGroups<T,Desc>>& particles
+ saw::data<sch::ParticleSpheroidGroup<T,Desc>>& particles
){
auto& info_f = fields.template get<"info">();
// Set everything as walls
@@ -103,15 +111,41 @@ saw::error_or<void> setup_initial_conditions(
// Particles
{
- saw::data<sch::Scalar<T>> rad_p, dense_p;
- rad_p.at({}).set(2.0);
+ saw::data<sch::Scalar<T>> dense_p;
dense_p.at({}).set(1);
- auto& spheroid_group = particles.template get<0u>();
- spheroid_group = create_spheroid_particle_group<T,Desc::D>(
- rad_p,
+ // auto& spheroid_group = particles.template get<0u>();
+ auto& spheroid_group = particles;
+
+ spheroid_group = create_spheroid_particle_group<T,Desc::D,2.0f>(
dense_p,
{64u}
);
+
+ {
+ auto& p = spheroid_group.template get<"particles">();
+
+ p = {{{16u}}};
+
+ iterator<1u>::apply(
+ [&](auto& index){
+ // Set Pos here?
+ auto& p_ind = p.at(index);
+
+ auto& p_rb = p_ind.template get<"rigid_body">();
+ auto& p_pos = p_rb.template get<"position">();
+
+ // TODO CONTINUE HERE NEED to init pos here !!!!
+
+ auto& p_pos_old = p_rb.template get<"position_old">();
+ p_pos_old = p_pos;
+ },
+ {},
+ p.meta()
+ );
+ }
+ }
+ // Particle in hacky flavour
+ {
}
return saw::make_void();
@@ -119,24 +153,104 @@ saw::error_or<void> setup_initial_conditions(
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::UInt64> t_i,
- device& dev
+ 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::Ptr<sch::ParticleSpheroidGroup<T,Desc>>,encode::Sycl<saw::encode::Native>>& p_group,
+ saw::data<sch::UInt64> t_i,
+ device& dev
){
auto& q = dev.get_handle();
auto& info_f = fields.template get<"info">();
- {
- }
+ auto& parts = p_group.template get<"particles">();
+ auto& p_mask = p_group.template get<"mask">();
+ auto& vels = macros.template get<"velocity">();
+ auto& forces = macros.template get<"force">();
+
+ auto p_meta = parts.meta();
+ q.submit([&](acpp::sycl::handler& h){
+
+ 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]);
+ }
+
+ // Reset the force to zero
+ forces.at(index) = {};
+ });
+ }).wait();
+
+ q.submit([&](acpp::sycl::handler& h){
+ h.parallel_for(acpp::sycl::range<1u>{p_meta.at({0u}).get()}, [=](acpp::sycl::id<1u> idx){
+
+ saw::data<sch::FixedArray<sch::UInt64,1u>> index;
+ for(uint64_t i = 0u; i < 1u; ++i){
+ index.at({{i}}).set(idx[i]);
+ }
+ auto& p = parts.at(index);
+ auto& p_rb = p.template get<"rigid_body">();
+ saw::data<sch::Scalar<T>> delta_t;
+ delta_t.at({}).set(1.0f);
+
+ auto& p_pos = p_rb.template get<"position">();
+ auto& p_rot = p_rb.template get<"rotation">();
+
+ iterator<Desc::D>::apply(
+ [&](auto& m_ind){
+ saw::data<sch::Vector<T,Desc::D>> index_shift;
+ for(uint64_t i{0u}; i < Desc::D; ++i){
+ index_shift.at({{i}}) = m_ind.at({i}).template cast_to<T>() - (vels.meta().at({i})+1u).template cast_to<T>() * 0.5;
+ }
+
+ saw::data<sch::Vector<T,Desc::D>> transformed_pos;
+ for(uint64_t i{0u}; i < Desc::D; ++i){
+ // TODO add rotation, scaling here.
+ transformed_pos.at({{i}}) = index_shift.at({{i}});
+ }
+
+ // Lagrange indicator position
+ auto p_pos_lag = p_pos + transformed_pos;
+
+ // Pick the closest velocity and clamp it
+ saw::data<sch::FixedArray<sch::UInt64,Desc::D>> p_cell_pos;
+ saw::data<sch::Vector<sch::UInt64,Desc::D>> p_cell_pos_vec;
+ for(uint64_t i{0u}; i < Desc::D; ++i){
+ p_cell_pos.at({{i}}) = (p_pos_lag.at({{i}}) + 0.5).template cast_to<sch::UInt64>();
+ p_cell_pos.at({{i}}).set(std::max(0ul,std::min(p_cell_pos.at({{i}}).get(), vels.meta().at({{i}}).get() - 2ul)));
+ p_cell_pos_vec.at({{i}}) = p_cell_pos.at({{i}});
+ }
+
+ auto& u_fluid = vels.at(p_cell_pos);
+
+ // this is our relative position to the particle
+ auto rel_cell_to_part_pos = p_cell_pos_vec.template cast_to<T>() - p_pos;
+
+ auto p_vel = (p_pos - p_rb.template get<"position_old">()) * delta_t;
+
+ auto u_solid = p_vel + saw::math::cross(p_rot,rel_cell_to_part_pos);
+
+ // Force
+ auto force = (u_solid - u_fluid) / delta_t;
+
+ // TODO HERE ATOMIC! !!!!
+ forces.at(p_cell_pos) = forces.at(p_cell_pos) + force;
+
+ // TODO APPLY FORCE TO PARTICLE
+ },
+ {},
+ p_mask.meta()
+ );
+
+ verlet_step_lambda<T,Desc::D>(p,delta_t);
+ });
+ }).wait();
// auto coll_ev =
q.submit([&](acpp::sycl::handler& h){
- saw::data<sch::Vector<T,Desc::D>> force;
- force.at({{1}}).set(-1.0);
// Need nicer things to handle the flow. I see improvement here
- component<T,Desc,cmpt::BGKGuo, encode::Sycl<saw::encode::Native>> collision{0.8};
+ component<T,Desc,cmpt::BGK, encode::Sycl<saw::encode::Native>> collision{0.8};
component<T,Desc,cmpt::BounceBack,encode::Sycl<saw::encode::Native>> bb;
h.parallel_for(acpp::sycl::range<Desc::D>{dim_x,dim_y}, [=](acpp::sycl::id<Desc::D> idx){
@@ -154,14 +268,12 @@ saw::error_or<void> step(
bb.apply(fields,index,t_i);
break;
case 2u:
- collision.apply(fields,macros,index,t_i);
+ // collision.apply(fields,macros,index,t_i);
break;
default:
break;
}
});
-
-
}).wait();
// Step
@@ -210,9 +322,11 @@ saw::error_or<void> lbm_main(int argc, char** argv){
// 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_particles_data = saw::data<sch::ParticleGroups<T,Desc>>();
+ auto lbm_particle_data_ptr = saw::heap<saw::data<sch::ParticleSpheroidGroup<T,Desc>>>();
+ // auto lbm_particles_ptr = saw::heap<saw::data<sch::FixedArray<sch::ParticleRigidBody<T,Desc::D>,part_count>>>();
+ // saw::data<sch::Array<T,Desc::D>> p_mask;
- std::cout<<"Estimated Bytes: "<<memory_estimate<sch::ChunkStruct<T,Desc>,sch::MacroStruct<T,Desc>>().get()<<std::endl;
+ std::cout<<"Estimated Bytes of LBM Fields: "<<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()){
@@ -236,7 +350,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,lbm_particles_data);
+ auto eov = setup_initial_conditions<T,Desc>(*lbm_data_ptr,*lbm_macro_data_ptr,*lbm_particle_data_ptr);
if(eov.is_error()){
return eov;
}
@@ -250,10 +364,8 @@ saw::error_or<void> lbm_main(int argc, char** argv){
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};
- sycl_q.wait();
+ saw::data<sch::ParticleSpheroidGroup<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_particle_data{sycl_q};
- 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()){
@@ -266,7 +378,24 @@ saw::error_or<void> 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();
+
+ auto lsd_view = make_view(lbm_sycl_data);
+ auto lsdm_view = make_view(lbm_sycl_macro_data);
+ auto lsdp_view = make_view(lbm_sycl_particle_data);
+
+ {
+ auto eov = write_vtk_file(out_dir,"ms",0u, *lbm_macro_data_ptr);
+ if(eov.is_error()){
+ return eov;
+ }
+ }
saw::data<sch::UInt64> time_steps{16u*4096ul};
auto& info_f = lsd_view.template get<"info">();
@@ -274,14 +403,13 @@ saw::error_or<void> lbm_main(int argc, char** argv){
for(saw::data<sch::UInt64> i{0u}; i < time_steps and krun; ++i){
// BC + Collision
{
- auto eov = step<T,Desc>(lsd_view,lsdm_view,i,dev);
+ auto eov = step<T,Desc>(lsd_view,lsdm_view,lsdp_view,i,dev);
if(eov.is_error()){
return eov;
}
}
sycl_q.wait();
- /*
- {
+ if(i.get()%1u ==0u){
{
auto eov = dev.copy_to_host(lbm_sycl_macro_data,*lbm_macro_data_ptr);
if(eov.is_error()){
@@ -295,7 +423,6 @@ saw::error_or<void> lbm_main(int argc, char** argv){
}
}
}
- */
/*
{
{
generated by cgit v1.2.3 (git 2.25.1) at 2026-05-11 16:18:13 +0000