diff options
author | Adrian Kummerlaender | 2021-05-17 00:15:33 +0200 |
---|---|---|
committer | Adrian Kummerlaender | 2021-05-17 00:15:33 +0200 |
commit | 4ec94c97879aafef15f7663135745e4ba61e62cf (patch) | |
tree | 322ae3f003892513f529842ff0b3fd100573b680 /tangle/LLBM/lattice.h | |
download | LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.gz LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.bz2 LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.lz LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.xz LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.tar.zst LiterateLB-4ec94c97879aafef15f7663135745e4ba61e62cf.zip |
Extract first public LiterateLB version
Diffstat (limited to 'tangle/LLBM/lattice.h')
-rw-r--r-- | tangle/LLBM/lattice.h | 131 |
1 files changed, 131 insertions, 0 deletions
diff --git a/tangle/LLBM/lattice.h b/tangle/LLBM/lattice.h new file mode 100644 index 0000000..8ba1d66 --- /dev/null +++ b/tangle/LLBM/lattice.h @@ -0,0 +1,131 @@ +#pragma once + +#include "memory.h" +#include "call_tag.h" +#include "operator.h" + +#include "propagate.h" +#include "kernel/initialize.h" +#include "kernel/executor.h" + +template <typename DESCRIPTOR, typename T, typename S=T> +class Lattice { +private: +const descriptor::Cuboid<DESCRIPTOR> _cuboid; + +CyclicPopulationBuffer<DESCRIPTOR,S> _population; + +public: +Lattice(descriptor::Cuboid<DESCRIPTOR> cuboid): + _cuboid(cuboid), + _population(cuboid) { } + +descriptor::Cuboid<DESCRIPTOR> cuboid() const { + return _cuboid; +} + +void stream() { + _population.stream(); +} + +template <typename... OPERATOR> +void apply(OPERATOR... ops) { + const auto block_size = 32; + const auto block_count = (_cuboid.volume + block_size - 1) / block_size; + kernel::call_operators<DESCRIPTOR,T,S,OPERATOR...><<<block_count,block_size>>>( + _population.view(), ops... + ); +} + +template <typename OPERATOR, typename... ARGS> +void apply(ARGS&&... args) { + call_operator<OPERATOR>(typename OPERATOR::call_tag{}, std::forward<ARGS&&>(args)...); +} + +template <typename OPERATOR, typename... ARGS> +void call_operator(tag::call_by_cell_id, DeviceBuffer<std::size_t>& cells, ARGS... args) { + const auto block_size = 32; + const auto block_count = (cells.size() + block_size - 1) / block_size; + kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>( + _population.view(), cells.device(), cells.size(), std::forward<ARGS>(args)... + ); +} + +template <typename OPERATOR, typename... ARGS> +void call_operator(tag::call_by_cell_id, DeviceBuffer<bool>& mask, ARGS... args) { + const auto block_size = 32; + const auto block_count = (_cuboid.volume + block_size - 1) / block_size; + kernel::call_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>( + _population.view(), mask.device(), std::forward<ARGS>(args)... + ); +} + +template <typename OPERATOR, typename... ARGS> +void call_operator(tag::call_by_list_index, std::size_t count, ARGS... args) { + const auto block_size = 32; + const auto block_count = (count + block_size - 1) / block_size; + kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>( + _population.view(), count, std::forward<ARGS>(args)... + ); +} + +template <typename FUNCTOR, typename... ARGS> +void inspect(ARGS&&... args) { + call_functor<FUNCTOR>(typename FUNCTOR::call_tag{}, std::forward<ARGS&&>(args)...); +} + +template <typename FUNCTOR, typename... ARGS> +void call_functor(tag::call_by_cell_id, DeviceBuffer<std::size_t>& cells, ARGS... args) { + const auto block_size = 32; + const auto block_count = (cells.size() + block_size - 1) / block_size; + kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>( + _population.view(), cells.device(), cells.size(), std::forward<ARGS>(args)... + ); +} + +template <typename FUNCTOR, typename... ARGS> +void call_functor(tag::call_by_cell_id, DeviceBuffer<bool>& mask, ARGS... args) { + const auto block_size = 32; + const auto block_count = (_cuboid.volume + block_size - 1) / block_size; + kernel::call_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>( + _population.view(), mask.device(), std::forward<ARGS>(args)... + ); +} + +template <typename FUNCTOR, typename... ARGS> +void call_functor(tag::call_by_spatial_cell_mask, DeviceBuffer<bool>& mask, ARGS... args) { + const dim3 block(32,8,4); + const dim3 grid((_cuboid.nX + block.x - 1) / block.x, + (_cuboid.nY + block.y - 1) / block.y, + (_cuboid.nZ + block.z - 1) / block.z); + kernel::call_spatial_functor<FUNCTOR,DESCRIPTOR,T,S,ARGS...><<<grid,block>>>( + _population.view(), mask.device(), std::forward<ARGS>(args)... + ); +} + +template <typename OPERATOR, typename... ARGS> +void helper(ARGS&&... args) { + tagged_helper<OPERATOR>(typename OPERATOR::call_tag{}, std::forward<ARGS&&>(args)...); +} + +template <typename OPERATOR, typename... ARGS> +void tagged_helper(tag::post_process_by_list_index, std::size_t count, ARGS... args) { + const auto block_size = 32; + const auto block_count = (count + block_size - 1) / block_size; + kernel::call_operator_using_list<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<block_count,block_size>>>( + DESCRIPTOR(), count, std::forward<ARGS>(args)... + ); +} + +template <typename OPERATOR, typename... ARGS> +void tagged_helper(tag::post_process_by_spatial_cell_mask, DeviceBuffer<bool>& mask, ARGS... args) { + const dim3 block(32,8,4); + const dim3 grid((_cuboid.nX + block.x - 1) / block.x, + (_cuboid.nY + block.y - 1) / block.y, + (_cuboid.nZ + block.z - 1) / block.z); + kernel::call_spatial_operator<OPERATOR,DESCRIPTOR,T,S,ARGS...><<<grid,block>>>( + _cuboid, mask.device(), std::forward<ARGS>(args)... + ); +} + +}; |