35 sham::DeviceBuffer<Tvec> pixel_to_positions(
36 Tvec center, Tvec delta_x, Tvec delta_y,
u32 nx,
u32 ny) {
38 sham::DeviceBuffer<Tvec> ret{nx * ny, shamsys::instance::get_compute_scheduler_ptr()};
40 sham::DeviceQueue &q = shamsys::instance::get_compute_scheduler().
get_queue();
43 q, sham::MultiRef{}, sham::MultiRef{ret}, nx * ny, [=](
u32 gid, Tvec *position) {
46 f64 fx = ((
f64(ix) + 0.5) / nx) - 0.5;
47 f64 fy = ((
f64(iy) + 0.5) / ny) - 0.5;
48 position[gid] = center + delta_x * fx + delta_y * fy;
55 sham::DeviceBuffer<shammath::Ray<Tvec>> pixel_to_orthographic_rays(
56 Tvec center, Tvec delta_x, Tvec delta_y,
u32 nx,
u32 ny) {
58 using Tscal = shambase::VecComponent<Tvec>;
60 sham::DeviceBuffer<shammath::Ray<Tvec>> ret{
61 nx * ny, shamsys::instance::get_compute_scheduler_ptr()};
63 sham::DeviceQueue &q = shamsys::instance::get_compute_scheduler().
get_queue();
65 Tvec e_z = sycl::cross(delta_x, delta_y);
66 Tscal len = sycl::length(e_z);
69 "The cross product of delta_x and delta_y is zero\n"
91 [=](
u32 gid, shammath::Ray<Tvec> *ray) {
94 f64 fx = ((
f64(ix) + 0.5) / nx) - 0.5;
95 f64 fy = ((
f64(iy) + 0.5) / ny) - 0.5;
96 Tvec pos_render = center + delta_x * fx + delta_y * fy;
98 ray[gid] = shammath::Ray<Tvec>(pos_render, e_z);
104 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
105 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_slice(
106 std::string field_name,
107 const sham::DeviceBuffer<Tvec> &positions,
108 std::optional<std::function<py::array_t<Tfield>(
size_t, shamrock::PatchDataLazyGetter &)>>
109 custom_getter) -> sham::DeviceBuffer<Tfield> {
113 "sph::CartesianRender",
115 "compute_slice field_name: {}, positions count: {}",
126 [&](
auto field_getter) -> sham::DeviceBuffer<Tfield> {
127 return compute_slice(field_getter, positions);
134 "sph::CartesianRender", sham::format(
"compute_slice took {}", t.
get_time_str()));
140 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
141 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_slice(
142 shamrock::solvergraph::Field<Tfield> &field,
const sham::DeviceBuffer<Tvec> &positions)
143 -> sham::DeviceBuffer<Tfield> {
145 if (field.get_nvar() != 1) {
147 "render only supports fields with nvar == 1");
150 shambase::DistributedData<u32>
sizes{};
151 scheduler().for_each_patchdata_nonempty(
152 [&](
const shamrock::patch::Patch p, shamrock::patch::PatchDataLayer &pdat) {
155 field.check_sizes(sizes);
158 = [&](
const shamrock::patch::Patch cur_p,
159 shamrock::patch::PatchDataLayer &pdat) ->
const sham::DeviceBuffer<Tfield> & {
160 return field.get_buf(cur_p.
id_patch);
163 return compute_slice(field_getter, positions);
166 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
167 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_column_integ(
168 std::string field_name,
169 const sham::DeviceBuffer<shammath::Ray<Tvec>> &rays,
170 std::optional<std::function<py::array_t<Tfield>(
size_t, shamrock::PatchDataLazyGetter &)>>
171 custom_getter) -> sham::DeviceBuffer<Tfield> {
175 "sph::CartesianRender",
177 "compute_column_integ field_name: {}, rays count: {}",
188 [&](
auto field_getter) -> sham::DeviceBuffer<Tfield> {
189 return compute_column_integ(field_getter, rays);
196 "sph::CartesianRender",
197 sham::format(
"compute_column_integ took {}", t.
get_time_str()));
203 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
204 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_column_integ(
205 shamrock::solvergraph::Field<Tfield> &field,
206 const sham::DeviceBuffer<shammath::Ray<Tvec>> &rays) -> sham::DeviceBuffer<Tfield> {
208 if (field.get_nvar() != 1) {
210 "render only supports fields with nvar == 1");
213 shambase::DistributedData<u32>
sizes{};
214 scheduler().for_each_patchdata_nonempty(
215 [&](
const shamrock::patch::Patch p, shamrock::patch::PatchDataLayer &pdat) {
218 field.check_sizes(sizes);
221 = [&](
const shamrock::patch::Patch cur_p,
222 shamrock::patch::PatchDataLayer &pdat) ->
const sham::DeviceBuffer<Tfield> & {
223 return field.get_buf(cur_p.
id_patch);
226 return compute_column_integ(field_getter, rays);
229 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
230 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_azymuthal_integ(
231 std::string field_name,
232 const sham::DeviceBuffer<shammath::RingRay<Tvec>> &ring_rays,
233 std::optional<std::function<py::array_t<Tfield>(
size_t, shamrock::PatchDataLazyGetter &)>>
234 custom_getter) -> sham::DeviceBuffer<Tfield> {
238 "sph::CartesianRender",
240 "compute_azymuthal_integ field_name: {}, ring_rays count: {}",
242 ring_rays.get_size()));
251 [&](
auto field_getter) -> sham::DeviceBuffer<Tfield> {
252 return compute_azymuthal_integ(field_getter, ring_rays);
259 "sph::CartesianRender",
260 sham::format(
"compute_azymuthal_integ took {}", t.
get_time_str()));
266 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
267 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_slice(
268 std::function<field_getter_t> field_getter,
const sham::DeviceBuffer<Tvec> &positions)
269 -> sham::DeviceBuffer<Tfield> {
271 auto part_counts = shamrock::solvergraph::Indexes<u32>::make_shared(
"part_counts",
"N");
273 = std::make_shared<shamrock::solvergraph::FieldRefs<Tvec>>(
"positions",
"\\mathbf{r}");
274 auto hpart_refs = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"h_part",
"h");
276 = std::make_shared<shamrock::solvergraph::Field<Tfield>>(1,
"field_data",
"f");
278 shamrock::solvergraph::DDPatchDataFieldRef<Tvec> pos_dd;
279 shamrock::solvergraph::DDPatchDataFieldRef<Tscal> h_dd;
281 scheduler().for_each_patchdata_nonempty(
282 [&](
const shamrock::patch::Patch cur_p, shamrock::patch::PatchDataLayer &pdat) {
287 pos_dd.add_obj(
id, std::ref(pdat.get_field<Tvec>(0)));
289 id, std::ref(pdat.get_field<Tscal>(pdat.pdl().
get_field_idx<Tscal>(
"hpart"))));
292 positions_refs->set_refs(pos_dd);
293 hpart_refs->set_refs(h_dd);
297 scheduler().for_each_patchdata_nonempty(
298 [&](
const shamrock::patch::Patch cur_p, shamrock::patch::PatchDataLayer &pdat) {
299 const sham::DeviceBuffer<Tfield> &src = field_getter(cur_p, pdat);
303 auto gpart_mass = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"gpart_mass",
"m");
304 gpart_mass->data = solver_config.gpart_mass;
306 auto tree_reduction_level
307 = shamrock::solvergraph::IDataEdge<u32>::make_shared(
"tree_reduction_level",
"l");
308 tree_reduction_level->data = solver_config.tree_reduction_level;
310 auto interp_points = std::make_shared<shamrock::solvergraph::DeviceBufferEdge<Tvec>>(
311 "interp_points",
"\\mathbf{q}");
312 interp_points->value.resize(positions.
get_size());
313 interp_points->value.copy_from(positions);
315 auto interpolated_field = std::make_shared<shamrock::solvergraph::DeviceBufferEdge<Tfield>>(
316 "interpolated_field",
"f_{\\rm interp}");
318 auto node = std::make_shared<SPHInterpolation<Tvec, Tfield, SPHKernel>>();
321 tree_reduction_level,
330 sham::DeviceBuffer<Tfield> ret{
331 interpolated_field->value.get_size(), shamsys::instance::get_compute_scheduler_ptr()};
332 ret.
copy_from(interpolated_field->value);
337 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
338 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_column_integ(
339 std::function<field_getter_t> field_getter,
340 const sham::DeviceBuffer<shammath::Ray<Tvec>> &rays) -> sham::DeviceBuffer<Tfield> {
342 auto part_counts = shamrock::solvergraph::Indexes<u32>::make_shared(
"part_counts",
"N");
344 = std::make_shared<shamrock::solvergraph::FieldRefs<Tvec>>(
"positions",
"\\mathbf{r}");
345 auto hpart_refs = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"h_part",
"h");
347 = std::make_shared<shamrock::solvergraph::Field<Tfield>>(1,
"field_data",
"f");
349 shamrock::solvergraph::DDPatchDataFieldRef<Tvec> pos_dd;
350 shamrock::solvergraph::DDPatchDataFieldRef<Tscal> h_dd;
352 scheduler().for_each_patchdata_nonempty(
353 [&](
const shamrock::patch::Patch cur_p, shamrock::patch::PatchDataLayer &pdat) {
358 pos_dd.add_obj(
id, std::ref(pdat.get_field<Tvec>(0)));
360 id, std::ref(pdat.get_field<Tscal>(pdat.pdl().
get_field_idx<Tscal>(
"hpart"))));
363 positions_refs->set_refs(pos_dd);
364 hpart_refs->set_refs(h_dd);
368 scheduler().for_each_patchdata_nonempty(
369 [&](
const shamrock::patch::Patch cur_p, shamrock::patch::PatchDataLayer &pdat) {
370 const sham::DeviceBuffer<Tfield> &src = field_getter(cur_p, pdat);
374 auto gpart_mass = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"gpart_mass",
"m");
375 gpart_mass->data = solver_config.gpart_mass;
377 auto tree_reduction_level
378 = shamrock::solvergraph::IDataEdge<u32>::make_shared(
"tree_reduction_level",
"l");
379 tree_reduction_level->data = solver_config.tree_reduction_level;
382 = std::make_shared<shamrock::solvergraph::DeviceBufferEdge<shammath::Ray<Tvec>>>(
383 "rays",
"\\mathbf{r}_{\\rm ray}");
384 rays_edge->value.resize(rays.get_size());
385 rays_edge->value.copy_from(rays);
387 auto interpolated_field = std::make_shared<shamrock::solvergraph::DeviceBufferEdge<Tfield>>(
388 "interpolated_field",
"f_{\\rm interp}");
390 auto node = std::make_shared<SPHColumnInteg<Tvec, Tfield, SPHKernel>>();
393 tree_reduction_level,
402 sham::DeviceBuffer<Tfield> ret{
403 interpolated_field->value.get_size(), shamsys::instance::get_compute_scheduler_ptr()};
404 ret.
copy_from(interpolated_field->value);
409 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
410 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_azymuthal_integ(
411 std::function<field_getter_t> field_getter,
412 const sham::DeviceBuffer<shammath::RingRay<Tvec>> &ring_rays)
413 -> sham::DeviceBuffer<Tfield> {
415 auto part_counts = shamrock::solvergraph::Indexes<u32>::make_shared(
"part_counts",
"N");
417 = std::make_shared<shamrock::solvergraph::FieldRefs<Tvec>>(
"positions",
"\\mathbf{r}");
418 auto hpart_refs = std::make_shared<shamrock::solvergraph::FieldRefs<Tscal>>(
"h_part",
"h");
420 = std::make_shared<shamrock::solvergraph::Field<Tfield>>(1,
"field_data",
"f");
422 shamrock::solvergraph::DDPatchDataFieldRef<Tvec> pos_dd;
423 shamrock::solvergraph::DDPatchDataFieldRef<Tscal> h_dd;
425 scheduler().for_each_patchdata_nonempty(
426 [&](
const shamrock::patch::Patch cur_p, shamrock::patch::PatchDataLayer &pdat) {
431 pos_dd.add_obj(
id, std::ref(pdat.get_field<Tvec>(0)));
433 id, std::ref(pdat.get_field<Tscal>(pdat.pdl().
get_field_idx<Tscal>(
"hpart"))));
436 positions_refs->set_refs(pos_dd);
437 hpart_refs->set_refs(h_dd);
441 scheduler().for_each_patchdata_nonempty(
442 [&](
const shamrock::patch::Patch cur_p, shamrock::patch::PatchDataLayer &pdat) {
443 const sham::DeviceBuffer<Tfield> &src = field_getter(cur_p, pdat);
447 auto gpart_mass = shamrock::solvergraph::IDataEdge<Tscal>::make_shared(
"gpart_mass",
"m");
448 gpart_mass->data = solver_config.gpart_mass;
450 auto tree_reduction_level
451 = shamrock::solvergraph::IDataEdge<u32>::make_shared(
"tree_reduction_level",
"l");
452 tree_reduction_level->data = solver_config.tree_reduction_level;
455 = std::make_shared<shamrock::solvergraph::DeviceBufferEdge<shammath::RingRay<Tvec>>>(
456 "ring_rays",
"\\mathbf{r}_{\\rm ring}");
457 ring_rays_edge->value.resize(ring_rays.get_size());
458 ring_rays_edge->value.copy_from(ring_rays);
460 auto interpolated_field = std::make_shared<shamrock::solvergraph::DeviceBufferEdge<Tfield>>(
461 "interpolated_field",
"f_{\\rm interp}");
463 auto node = std::make_shared<SPHAzymuthalInteg<Tvec, Tfield, SPHKernel>>();
466 tree_reduction_level,
475 sham::DeviceBuffer<Tfield> ret{
476 interpolated_field->value.get_size(), shamsys::instance::get_compute_scheduler_ptr()};
477 ret.
copy_from(interpolated_field->value);
482 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
483 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_azymuthal_integ(
484 shamrock::solvergraph::Field<Tfield> &field,
485 const sham::DeviceBuffer<shammath::RingRay<Tvec>> &ring_rays)
486 -> sham::DeviceBuffer<Tfield> {
488 if (field.get_nvar() != 1) {
490 "render only supports fields with nvar == 1");
493 shambase::DistributedData<u32>
sizes{};
494 scheduler().for_each_patchdata_nonempty(
495 [&](
const shamrock::patch::Patch p, shamrock::patch::PatchDataLayer &pdat) {
498 field.check_sizes(sizes);
501 = [&](
const shamrock::patch::Patch cur_p,
502 shamrock::patch::PatchDataLayer &pdat) ->
const sham::DeviceBuffer<Tfield> & {
503 return field.get_buf(cur_p.
id_patch);
506 return compute_azymuthal_integ(field_getter, ring_rays);
509 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
510 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_slice(
511 std::function<field_getter_t> field_getter,
516 u32 ny) -> sham::DeviceBuffer<Tfield> {
518 auto positions = pixel_to_positions(center, delta_x, delta_y, nx, ny);
520 return compute_slice(field_getter, positions);
523 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
524 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_column_integ(
525 std::function<field_getter_t> field_getter,
530 u32 ny) -> sham::DeviceBuffer<Tfield> {
532 auto rays = pixel_to_orthographic_rays(center, delta_x, delta_y, nx, ny);
534 return compute_column_integ(field_getter, rays);
537 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
538 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_slice(
539 shamrock::solvergraph::Field<Tfield> &field,
544 u32 ny) -> sham::DeviceBuffer<Tfield> {
546 auto positions = pixel_to_positions(center, delta_x, delta_y, nx, ny);
548 return compute_slice(field, positions);
551 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
552 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_column_integ(
553 shamrock::solvergraph::Field<Tfield> &field,
558 u32 ny) -> sham::DeviceBuffer<Tfield> {
560 auto rays = pixel_to_orthographic_rays(center, delta_x, delta_y, nx, ny);
562 return compute_column_integ(field, rays);
565 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
566 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_slice(
567 std::string field_name,
574 std::function<pybind11::array_t<Tfield>(
size_t, shamrock::PatchDataLazyGetter &)>>
575 custom_getter) -> sham::DeviceBuffer<Tfield> {
576 auto positions = pixel_to_positions(center, delta_x, delta_y, nx, ny);
577 return compute_slice(std::move(field_name), positions, std::move(custom_getter));
580 template<
class Tvec,
class Tfield,
template<
class>
class SPHKernel>
581 auto CartesianRender<Tvec, Tfield, SPHKernel>::compute_column_integ(
582 std::string field_name,
589 std::function<pybind11::array_t<Tfield>(
size_t, shamrock::PatchDataLazyGetter &)>>
590 custom_getter) -> sham::DeviceBuffer<Tfield> {
591 auto rays = pixel_to_orthographic_rays(center, delta_x, delta_y, nx, ny);
592 return compute_column_integ(std::move(field_name), rays, std::move(custom_getter));
Solver graph edge wrapping a global device buffer.
constexpr const char * sizes
Temporary sizes for h-iteration.
constexpr const char * part_counts
Particle counts per patch.
SPH azimuthal integration solver graph node.
SPH column integration solver graph node.
SPH slice interpolation solver graph node.
double f64
Alias for double.
std::uint32_t u32
32 bit unsigned integer
std::uint64_t u64
64 bit unsigned integer
void copy_from(const DeviceBuffer< T, new_target > &other, size_t copy_size)
Copies the content of another buffer to this one.
size_t get_size() const
Gets the number of elements in the buffer.
DeviceQueue & get_queue(u32 id=0)
Get a reference to a DeviceQueue.
iterator add_obj(u64 id, T &&obj)
Adds a new object to the collection.
std::string get_time_str() const
Converts the stored nanosecond time to a string representation.
void start()
Starts the timer.
void stop()
Stops the timer and stores the elapsed time in nanoseconds.
u32 get_field_idx(const std::string &field_name) const
Get the field id if matching name & type.
u32 get_obj_cnt() const
get the number of objects (particles) stored in this layer
This header file contains utility functions related to exception handling in the code.
void kernel_call(sham::DeviceQueue &q, RefIn in, RefOut in_out, u32 n, Functor &&func, SourceLocation &&callsite=SourceLocation{})
Submit a kernel to a SYCL queue.
ExcptTypes make_except_with_loc(std::string message, SourceLocation loc=SourceLocation{})
Create an exception with a message and a location.
i32 world_rank()
Gives the rank of the current process in the MPI communicator.
namespace for math utility
namespace for the sph model modules
u64 id_patch
unique key that identify the patch