summaryrefslogtreecommitdiff
path: root/lib
diff options
context:
space:
mode:
Diffstat (limited to 'lib')
-rw-r--r--lib/core/c++/collision.hpp34
-rw-r--r--lib/core/c++/macroscopic.hpp25
-rw-r--r--lib/sycl/c++/data.hpp48
3 files changed, 54 insertions, 53 deletions
diff --git a/lib/core/c++/collision.hpp b/lib/core/c++/collision.hpp
index 2f20601..85c3af9 100644
--- a/lib/core/c++/collision.hpp
+++ b/lib/core/c++/collision.hpp
@@ -45,43 +45,19 @@ public:
* Raw setup
*/
template<typename CellFieldSchema>
- void apply(saw::data<CellFieldSchema, Encode>& field, saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> index, saw::data<sch::UInt64> time_step){
+ void apply(const saw::data<CellFieldSchema, Encode>& field, saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>> index, saw::data<sch::UInt64> time_step) const {
bool is_even = ((time_step.get() % 2) == 0);
- auto& cell = field.at(index);
- auto& dfs_old = (is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
- // auto& dfs = (not is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
+ // auto& dfs_f = (is_even) ? field.template get<"dfs">() : field.template get<"dfs_old">();
+ auto& dfs_old_f = (is_even) ? field.template get<"dfs_old">() : field.template get<"dfs">();
saw::data<T> rho;
saw::data<sch::FixedArray<T,Descriptor::D>> vel;
- compute_rho_u<T,Descriptor>(dfs_old,rho,vel);
- auto eq = equilibrium<T,Descriptor>(rho,vel);
-
- for(uint64_t i = 0u; i < Descriptor::Q; ++i){
- dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i}));
- // dfs_old({i}).set(dfs_old({i}).get() + (1.0 / relaxation_) * (eq.at(i).get() - dfs_old({i}).get()));
- }
- }
-
- template<typename CellStructSchema>
- void apply(
- saw::data<CellStructSchema, Encode>* field,
- const saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>>& meta,
- const saw::data<sch::FixedArray<sch::UInt64,Descriptor::D>>& index,
- saw::data<sch::UInt64> time_step
- ){
- bool is_even = ((time_step.get() % 2) == 0);
- auto& cell = field[0u];
-
- auto& dfs_old = (is_even) ? cell.template get<"dfs_old">() : cell.template get<"dfs">();
-
- saw::data<T> rho;
- saw::data<sch::FixedArray<T,Descriptor::D>> vel;
- compute_rho_u<T,Descriptor>(dfs_old,rho,vel);
+ compute_rho_u<T,Descriptor>(dfs_old_f.at(index),rho,vel);
auto eq = equilibrium<T,Descriptor>(rho,vel);
for(uint64_t i = 0u; i < Descriptor::Q; ++i){
- dfs_old({i}) = dfs_old({i}) + frequency_ * (eq.at(i) - dfs_old({i}));
+ dfs_old_f.at(index).at({i}) = dfs_old_f.at(index).at({i}) + frequency_ * (eq.at(i) - dfs_old_f.at(index).at({i}));
}
}
};
diff --git a/lib/core/c++/macroscopic.hpp b/lib/core/c++/macroscopic.hpp
index 51368e9..1532873 100644
--- a/lib/core/c++/macroscopic.hpp
+++ b/lib/core/c++/macroscopic.hpp
@@ -4,6 +4,31 @@
namespace kel {
namespace lbm {
+template<typename T, typename Desc>
+void compute_rho_u (
+ const saw::data<sch::FixedArray<T,Desc::Q>>& dfs,
+ saw::data<T>& rho,
+ saw::data<sch::FixedArray<T,Desc::D>>& vel
+ )
+{
+ using dfi = df_info<T, Desc>;
+
+ rho = 0;
+ for(size_t j = 0; j < Desc::D; ++j){
+ vel.at({{j}}).set(0);
+ }
+
+ for(size_t j = 0; j < Desc::Q; ++j){
+ rho = rho + dfs.at({j});
+ for(size_t i = 0; i < Desc::D; ++i){
+ vel.at({{i}}) = vel.at({{i}}) + dfi::directions[j][i] * dfs.at({j}).get();
+ }
+ }
+
+ for(size_t i = 0; i < Desc::D; ++i){
+ vel.at({i}) = vel.at({i}) / rho;
+ }
+}
/**
* Calculate the macroscopic variables rho and u in Lattice Units.
*/
diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp
index cffeb38..d5adb95 100644
--- a/lib/sycl/c++/data.hpp
+++ b/lib/sycl/c++/data.hpp
@@ -25,15 +25,14 @@ public:
using Schema = schema::FixedArray<Sch,Dims...>;
private:
acpp::sycl::queue* q_;
- data<Sch>* values_;
+ data<Sch,Encode>* values_;
- SAW_FORBID_COPY(data);
+ //SAW_FORBID_COPY(data);
public:
data(acpp::sycl::queue& q__):
q_{&q__},
values_{nullptr}
{
- std::cout<<"Hey: "<<ct_multiply<uint64_t,Dims...>::value<<std::endl;
values_ = acpp::sycl::malloc_device<data<Sch,Encode>>(ct_multiply<uint64_t,Dims...>::value,*q_);
SAW_ASSERT(values_ and q_);
}
@@ -49,14 +48,18 @@ public:
}
static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() {
- return {std::array<uint64_t, sizeof...(Dims)>{Dims...}};
+ return saw::data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>>{{Dims...}};
}
- data<Sch>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index){
+ constexpr data<Sch,Encode>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index){
return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
}
- constexpr data<Sch,encode::Native>* flat_data() {
+ constexpr data<Sch,Encode>& at(data<schema::FixedArray<schema::UInt64,sizeof...(Dims)>>& index) const{
+ return values_[kel::lbm::flatten_index<schema::UInt64,sizeof...(Dims)>::apply(index,get_dims()).get()];
+ }
+
+ constexpr data<Sch,Encode>* flat_data() {
return values_;
}
};
@@ -79,7 +82,7 @@ public:
return values_.at(index);
}
- const data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+ data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& ghost_at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
return values_.at(index);
}
@@ -87,7 +90,7 @@ public:
return data<InnerSchema,kel::lbm::encode::Sycl<Encode>>::get_dims();
}
- data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
+ data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index){
std::decay_t<decltype(index)> ind;
for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
ind.at({i}) = index.at({i}) + Ghost;
@@ -95,7 +98,7 @@ public:
return values_.at(ind);
}
- const data<ValueSchema, kel::lbm::encode::Sycl<Encode>>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
+ data<ValueSchema, Encode>& at(const data<schema::FixedArray<schema::UInt64,sizeof...(Sides)>>& index) const {
std::decay_t<decltype(index)> ind;
for(uint64_t i = 0u; i < sizeof...(Sides); ++i){
ind.at({i}) = index.at({i}) + Ghost;
@@ -116,12 +119,11 @@ public:
}
};
-template<uint64_t Ghost, uint64_t... Meta, typename... Sch, string_literal... Keys, typename Encode>
-class data<schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>, kel::lbm::encode::Sycl<Encode> > final {
+template<typename... Members, typename Encode>
+class data<schema::Struct<Members...>, kel::lbm::encode::Sycl<Encode> > final {
public:
- static constexpr data<schema::FixedArray<schema::UInt64,sizeof...(Meta)>> meta = {{Meta...}};
- using StorageT = std::tuple<data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>...>;
- using Schema = schema::Struct<schema::Member<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>, Keys>...>;
+ using StorageT = std::tuple<data<typename Members::ValueType,kel::lbm::encode::Sycl<Encode>>...>;
+ using Schema = schema::Struct<Members...>;
private:
/**
@@ -133,24 +135,24 @@ private:
template<std::size_t... Is>
constexpr data(acpp::sycl::queue& q, std::index_sequence<Is...>):
members_{(static_cast<void>(Is), q)...}
- {
-
- }
+ {}
public:
data(acpp::sycl::queue& q__):
- data{q__, std::make_index_sequence<sizeof...(Sch)>{}}
-// members_(q__)//data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>(q__)...)
+ data{q__, std::make_index_sequence<sizeof...(Members)>{}}
{
q__.wait();
}
template<saw::string_literal K>
auto& get(){
- return std::get<parameter_key_pack_index<K, Keys...>::value>(members_);
+ return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_);
+ }
+
+ template<saw::string_literal K>
+ const auto& get() const {
+ return std::get<parameter_key_pack_index<K, Members::KeyLiteral...>::value>(members_);
}
};
-
-
}
namespace kel {
@@ -175,8 +177,6 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final {
static_assert(sizeof(std::decay_t<decltype(sycl_ptr)>) == sizeof(std::decay_t<decltype(host_ptr)>), "Unequal size");
- std::cout<<host_member_data.flat_size().get()<<" "<<std::endl;
-
q.submit([&](acpp::sycl::handler& h){
h.copy(host_ptr,sycl_ptr, host_member_data.flat_size().get());
}).wait();
generated by cgit v1.2.3 (git 2.25.1) at 2026-01-27 22:26:24 +0000