diff options
| -rw-r--r-- | .nix/badchimp.nix | 12 | ||||
| -rw-r--r-- | examples/poiseulle_particles_2d_gpu/sim.cpp | 8 | ||||
| -rw-r--r-- | lib/core/c++/chunk.hpp | 2 | ||||
| -rw-r--r-- | lib/sycl/c++/data.hpp | 72 | ||||
| -rw-r--r-- | vtk_to_png.py | 75 |
5 files changed, 109 insertions, 60 deletions
diff --git a/.nix/badchimp.nix b/.nix/badchimp.nix new file mode 100644 index 0000000..f747ffe --- /dev/null +++ b/.nix/badchimp.nix @@ -0,0 +1,12 @@ +{ lib +, stdenv +, fetchFromGitHub +}: + +stdenv.mkDerivation { + pname = "BadChimp"; + version = "0.0.0"; + + src = fetchFromGitHub { + }; +} diff --git a/examples/poiseulle_particles_2d_gpu/sim.cpp b/examples/poiseulle_particles_2d_gpu/sim.cpp index 9646b0a..ad60813 100644 --- a/examples/poiseulle_particles_2d_gpu/sim.cpp +++ b/examples/poiseulle_particles_2d_gpu/sim.cpp @@ -140,13 +140,7 @@ saw::error_or<void> lbm_main(int argc, char** argv){ } } - saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data; - { - auto eov = dev.allocate_on_device(lbm_sycl_data); - if(eov.is_error()){ - return eov; - } - } + saw::data<sch::ChunkStruct<T,Desc>, encode::Sycl<saw::encode::Native>> lbm_sycl_data{sycl_q}; { auto eov = dev.copy_to_device(lbm_data,lbm_sycl_data); if(eov.is_error()){ diff --git a/lib/core/c++/chunk.hpp b/lib/core/c++/chunk.hpp index 6b028ba..1da2c9f 100644 --- a/lib/core/c++/chunk.hpp +++ b/lib/core/c++/chunk.hpp @@ -81,7 +81,7 @@ public: } static constexpr auto flat_size() { - return data<InnerSchema,Encode>::flat_size(); + return data<InnerSchema,Encode>::flat_dims(); } }; diff --git a/lib/sycl/c++/data.hpp b/lib/sycl/c++/data.hpp index d9976dd..c6ea281 100644 --- a/lib/sycl/c++/data.hpp +++ b/lib/sycl/c++/data.hpp @@ -1,5 +1,3 @@ -#pragma once - #include "common.hpp" #include <kel/lbm/lbm.hpp> @@ -35,7 +33,7 @@ public: q_{&q__}, values_{nullptr} { - values_ = acpp::sycl::malloc<data<Sch>>(ct_multiply<uint64_t,Dims...>::value,q_); + values_ = acpp::sycl::malloc_device<data<Sch>>(ct_multiply<uint64_t,Dims...>::value,*q_); } ~data(){ @@ -43,7 +41,7 @@ public: return; } - acpp::sycl::free(values_,q_); + acpp::sycl::free(values_,*q_); } static constexpr data<schema::FixedArray<schema::UInt64, sizeof...(Dims)>> get_dims() { @@ -127,10 +125,20 @@ private: * Do it here by specializing. */ StorageT members_; + + 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__): - members_{{data<kel::lbm::sch::Chunk<Sch,Ghost,Meta...>,kel::lbm::encode::Sycl<Encode>>{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__)...) + { + q__.wait(); + } template<saw::string_literal K> auto& get(){ @@ -145,37 +153,6 @@ namespace kel { namespace lbm { namespace impl { template<typename Sch, typename Encode> -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 saw::data<Schema,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>(1u,q); - - return allocate_on_device_member<i+1u>(storage,q); - } - - return saw::make_void(); - } - - 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}; - - return eov; - } -}; - -template<typename Sch, typename Encode> struct sycl_copy_helper; template<typename... Members, typename Encode> @@ -186,13 +163,13 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final { static saw::error_or<void> copy_to_device_member(saw::data<Schema,Encode>& host_data, saw::data<Schema,encode::Sycl<Encode>>& sycl_data, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ using M = typename saw::parameter_pack_type<i,Members...>::type; - auto& host_member_data = host_data.template get<i>(); - auto& sycl_member_data = sycl_data.template get<i>(); + auto& host_member_data = host_data.template get<M::KeyLiteral>(); + auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>(); auto host_ptr = host_member_data.flat_data(); auto sycl_ptr = sycl_member_data.flat_data(); - q.memcpy(host_ptr, sycl_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size() ); + q.memcpy(host_ptr, sycl_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size().get() ); return copy_to_device_member<i+1u>(host_data,sycl_data,q); } @@ -209,13 +186,13 @@ struct sycl_copy_helper<sch::Struct<Members...>, Encode> final { static saw::error_or<void> copy_to_host_member(saw::data<Schema,encode::Sycl<Encode>>& sycl_data, saw::data<Schema,Encode>& host_data, sycl::queue& q){ if constexpr (i < sizeof...(Members)){ using M = typename saw::parameter_pack_type<i,Members...>::type; - auto& host_member_data = host_data.template get<i>(); - auto& sycl_member_data = sycl_data.template get<i>(); + auto& host_member_data = host_data.template get<M::KeyLiteral>(); + auto& sycl_member_data = sycl_data.template get<M::KeyLiteral>(); auto host_ptr = host_member_data.flat_data(); auto sycl_ptr = sycl_member_data.flat_data(); - q.memcpy(sycl_ptr, host_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size() ); + q.memcpy(sycl_ptr, host_ptr, sizeof(std::decay_t<decltype(host_ptr)>) * host_member_data.flat_size().get() ); return copy_to_host_member<i+1u>(sycl_data,host_data,q); } @@ -240,15 +217,6 @@ public: ~device() = default; template<typename Sch, typename Encode> - 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; - } - return saw::make_void(); - } - - template<typename Sch, typename Encode> saw::error_or<void> copy_to_device(saw::data<Sch,Encode>& host_data, saw::data<Sch,encode::Sycl<Encode>>& sycl_data){ return impl::sycl_copy_helper<Sch,Encode>::copy_to_device(host_data, sycl_data, q_); } diff --git a/vtk_to_png.py b/vtk_to_png.py new file mode 100644 index 0000000..9c3f4f3 --- /dev/null +++ b/vtk_to_png.py @@ -0,0 +1,75 @@ +import vtk +import glob +import os + +vtk_files = sorted(glob.glob("*.vtk")) +output_dir = "frames" +os.makedirs(output_dir, exist_ok=True) + +renderer = vtk.vtkRenderer() +render_window = vtk.vtkRenderWindow() +render_window.SetOffScreenRendering(1) +render_window.SetSize(1024, 512) +render_window.AddRenderer(renderer) + +renderer.SetBackground(1, 1, 1) + +camera_initialized = False + +for i, vtk_file in enumerate(vtk_files): + reader = vtk.vtkStructuredPointsReader() + reader.SetFileName(vtk_file) + reader.Update() + + image = reader.GetOutput() + + mag = vtk.vtkImageMagnitude(); + mag.SetInputData(image); + mag.Update(); + image = mag.GetOutput(); + scalar_name = "Velocity magnitude"; + + scalars = image.GetPointData().GetScalars() + if scalars is None: + raise RuntimeError("No scalar data found") + + vmin, vmax = scalars.GetRange() + + # Image slice mapper (2D!) + slice_mapper = vtk.vtkImageSliceMapper() + slice_mapper.SetInputData(image) + + # Pick the only slice (or middle if dimension >1) + extent = image.GetExtent() + slice_mapper.SetSliceNumber(extent[4]) # Z slice + + slice_actor = vtk.vtkImageSlice() + slice_actor.SetMapper(slice_mapper) + + # Grayscale mapping + prop = slice_actor.GetProperty() + prop.SetColorWindow(vmax - vmin) + prop.SetColorLevel(0.5 * (vmax + vmin)) + + renderer.RemoveAllViewProps() + renderer.AddViewProp(slice_actor) + + if not camera_initialized: + renderer.ResetCamera() + camera_initialized = True + + render_window.Render() + + # Capture to PNG + w2i = vtk.vtkWindowToImageFilter() + w2i.SetInput(render_window) + w2i.ReadFrontBufferOff() + w2i.Update() + + writer = vtk.vtkPNGWriter() + writer.SetFileName(f"{output_dir}/frame_{i:04d}.png") + writer.SetInputConnection(w2i.GetOutputPort()) + writer.Write() + + print(f"Saved frame {i} range=({vmin},{vmax})") + |