reduction_alg.cpp File Reference

#include "shambase/exception.hpp"
#include "shambase/string.hpp"
#include "shamalgs/memory.hpp"
#include "shamalgs/numeric.hpp"
#include "shambackends/DeviceQueue.hpp"
#include "shambackends/DeviceScheduler.hpp"
#include "shambackends/kernel_call.hpp"
#include "shambackends/math.hpp"
#include "shambackends/sycl.hpp"
#include "shamtree/kernels/reduction_alg.hpp"
#include <algorithm>
#include <memory>
#include <vector>

Include dependency graph for reduction_alg.cpp:

Go to the source code of this file.

Functions
template<class u_morton , class kername , class split_int >
void	sycl_generate_split_table (sycl::queue &queue, u32 morton_count, std::unique_ptr< sycl::buffer< u_morton > > &buf_morton, std::unique_ptr< sycl::buffer< split_int > > &buf_split_table)
template<class u_morton , class split_int >
void	sycl_generate_split_table (sham::DeviceQueue &queue, u32 morton_count, sham::DeviceBuffer< u_morton > &buf_morton, sham::DeviceBuffer< split_int > &buf_split_table)
template<class u_morton , class kername , class split_int >
void	sycl_reduction_iteration (sycl::queue &queue, u32 morton_count, std::unique_ptr< sycl::buffer< u_morton > > &buf_morton, std::unique_ptr< sycl::buffer< split_int > > &buf_split_table_in, std::unique_ptr< sycl::buffer< split_int > > &buf_split_table_out)
template<class u_morton , class split_int >
void	sycl_reduction_iteration (sham::DeviceQueue &queue, u32 morton_count, sham::DeviceBuffer< u_morton > &buf_morton, sham::DeviceBuffer< split_int > &buf_split_table_in, sham::DeviceBuffer< split_int > &buf_split_table_out)
void	update_morton_buf (sycl::queue &queue, u32 len, u32 val_ins, sycl::buffer< u32 > &buf_src, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map)
void	update_morton_buf (sham::DeviceQueue &queue, u32 len, u32 val_ins, sham::DeviceBuffer< u32 > &buf_src, sham::DeviceBuffer< u32 > &buf_reduc_index_map)
template<class split_int >
void	make_indexmap (sycl::queue &queue, u32 morton_count, u32 &morton_leaf_count, std::unique_ptr< sycl::buffer< split_int > > &buf_split_table, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map)
template<class split_int >
reduc_ret_t< split_int >	make_indexmap (const sham::DeviceScheduler_ptr &dev_sched, u32 morton_count, sham::DeviceBuffer< split_int > &buf_split_table)
template<class u_morton , class kername_split , class kername_reduc_it >
void	reduction_alg_impl (sycl::queue &queue, u32 morton_count, std::unique_ptr< sycl::buffer< u_morton > > &buf_morton, u32 reduction_level, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, u32 &morton_leaf_count)
template<class u_morton >
reduc_ret_t< u32 >	reduction_alg_impl (const sham::DeviceScheduler_ptr &dev_sched, u32 morton_count, sham::DeviceBuffer< u_morton > &buf_morton, u32 reduction_level)
template<>
void	reduction_alg< u32 > (sycl::queue &queue, u32 morton_count, std::unique_ptr< sycl::buffer< u32 > > &buf_morton, u32 reduction_level, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, u32 &morton_leaf_count)
template<>
void	reduction_alg< u64 > (sycl::queue &queue, u32 morton_count, std::unique_ptr< sycl::buffer< u64 > > &buf_morton, u32 reduction_level, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, u32 &morton_leaf_count)
template<>
reduc_ret_t< u32 >	reduction_alg< u32 > (const sham::DeviceScheduler_ptr &dev_sched, u32 morton_count, sham::DeviceBuffer< u32 > &buf_morton, u32 reduction_level)
template<>
reduc_ret_t< u32 >	reduction_alg< u64 > (const sham::DeviceScheduler_ptr &dev_sched, u32 morton_count, sham::DeviceBuffer< u64 > &buf_morton, u32 reduction_level)
template<class u_morton , class kername >
void	__sycl_morton_remap_reduction (sycl::queue &queue, u32 morton_leaf_count, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, std::unique_ptr< sycl::buffer< u_morton > > &buf_morton, std::unique_ptr< sycl::buffer< u_morton > > &buf_leaf_morton)
template<>
void	sycl_morton_remap_reduction< u32 > (sycl::queue &queue, u32 morton_leaf_count, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, std::unique_ptr< sycl::buffer< u32 > > &buf_morton, std::unique_ptr< sycl::buffer< u32 > > &buf_leaf_morton)
template<>
void	sycl_morton_remap_reduction< u64 > (sycl::queue &queue, u32 morton_leaf_count, std::unique_ptr< sycl::buffer< u32 > > &buf_reduc_index_map, std::unique_ptr< sycl::buffer< u64 > > &buf_morton, std::unique_ptr< sycl::buffer< u64 > > &buf_leaf_morton)
template<class u_morton >
void	sycl_morton_remap_reduction (sham::DeviceQueue &queue, u32 morton_leaf_count, sham::DeviceBuffer< u32 > &buf_reduc_index_map, sham::DeviceBuffer< u_morton > &buf_morton, sham::DeviceBuffer< u_morton > &buf_leaf_morton)
	Remaps a Morton tree on device using a reduction index map.
template void	sycl_morton_remap_reduction< u32 > (sham::DeviceQueue &queue, u32 morton_leaf_count, sham::DeviceBuffer< u32 > &buf_reduc_index_map, sham::DeviceBuffer< u32 > &buf_morton, sham::DeviceBuffer< u32 > &buf_leaf_morton)
template void	sycl_morton_remap_reduction< u64 > (sham::DeviceQueue &queue, u32 morton_leaf_count, sham::DeviceBuffer< u32 > &buf_reduc_index_map, sham::DeviceBuffer< u64 > &buf_morton, sham::DeviceBuffer< u64 > &buf_leaf_morton)

Detailed Description

Author

Timothée David–Cléris (tim.s.nosp@m.hamr.nosp@m.ock@p.nosp@m.roto.nosp@m.n.me)

Definition in file reduction_alg.cpp.

Macro Definition Documentation

NEW_BEHAVIOR

#define NEW_BEHAVIOR

Definition at line 133 of file reduction_alg.cpp.

OFFSET

#define OFFSET

Definition at line 136 of file reduction_alg.cpp.

Function Documentation

__sycl_morton_remap_reduction()

template<class u_morton , class kername >

void __sycl_morton_remap_reduction	(	sycl::queue &	queue,
		u32	morton_leaf_count,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map,
		std::unique_ptr< sycl::buffer< u_morton > > &	buf_morton,
		std::unique_ptr< sycl::buffer< u_morton > > &	buf_leaf_morton
	)

Definition at line 514 of file reduction_alg.cpp.

make_indexmap() [1/2]

template<class split_int >

reduc_ret_t< split_int > make_indexmap	(	const sham::DeviceScheduler_ptr &	dev_sched,
		u32	morton_count,
		sham::DeviceBuffer< split_int > &	buf_split_table
	)

Definition at line 358 of file reduction_alg.cpp.

make_indexmap() [2/2]

template<class split_int >

void make_indexmap	(	sycl::queue &	queue,
		u32	morton_count,
		u32 &	morton_leaf_count,
		std::unique_ptr< sycl::buffer< split_int > > &	buf_split_table,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map
	)

Definition at line 298 of file reduction_alg.cpp.

reduction_alg< u32 >() [1/2]

template<>

reduc_ret_t< u32 > reduction_alg< u32 >	(	const sham::DeviceScheduler_ptr &	dev_sched,
		u32	morton_count,
		sham::DeviceBuffer< u32 > &	buf_morton,
		u32	reduction_level
	)

Definition at line 476 of file reduction_alg.cpp.

reduction_alg< u32 >() [2/2]

template<>

void reduction_alg< u32 >	(	sycl::queue &	queue,
		u32	morton_count,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_morton,
		u32	reduction_level,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map,
		u32 &	morton_leaf_count
	)

Definition at line 459 of file reduction_alg.cpp.

reduction_alg< u64 >() [1/2]

template<>

reduc_ret_t< u32 > reduction_alg< u64 >	(	const sham::DeviceScheduler_ptr &	dev_sched,
		u32	morton_count,
		sham::DeviceBuffer< u64 > &	buf_morton,
		u32	reduction_level
	)

Definition at line 476 of file reduction_alg.cpp.

reduction_alg< u64 >() [2/2]

template<>

void reduction_alg< u64 >	(	sycl::queue &	queue,
		u32	morton_count,
		std::unique_ptr< sycl::buffer< u64 > > &	buf_morton,
		u32	reduction_level,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map,
		u32 &	morton_leaf_count
	)

Definition at line 476 of file reduction_alg.cpp.

reduction_alg_impl() [1/2]

template<class u_morton >

reduc_ret_t< u32 > reduction_alg_impl	(	const sham::DeviceScheduler_ptr &	dev_sched,
		u32	morton_count,
		sham::DeviceBuffer< u_morton > &	buf_morton,
		u32	reduction_level
	)

Definition at line 414 of file reduction_alg.cpp.

reduction_alg_impl() [2/2]

template<class u_morton , class kername_split , class kername_reduc_it >

void reduction_alg_impl	(	sycl::queue &	queue,
		u32	morton_count,
		std::unique_ptr< sycl::buffer< u_morton > > &	buf_morton,
		u32	reduction_level,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map,
		u32 &	morton_leaf_count
	)

Definition at line 376 of file reduction_alg.cpp.

sycl_generate_split_table() [1/2]

template<class u_morton , class split_int >

void sycl_generate_split_table	(	sham::DeviceQueue &	queue,
		u32	morton_count,
		sham::DeviceBuffer< u_morton > &	buf_morton,
		sham::DeviceBuffer< split_int > &	buf_split_table
	)

Definition at line 63 of file reduction_alg.cpp.

sycl_generate_split_table() [2/2]

template<class u_morton , class kername , class split_int >

void sycl_generate_split_table	(	sycl::queue &	queue,
		u32	morton_count,
		std::unique_ptr< sycl::buffer< u_morton > > &	buf_morton,
		std::unique_ptr< sycl::buffer< split_int > > &	buf_split_table
	)

Definition at line 34 of file reduction_alg.cpp.

sycl_morton_remap_reduction()

template<class u_morton >

void sycl_morton_remap_reduction	(	sham::DeviceQueue &	queue,
		u32	morton_leaf_count,
		sham::DeviceBuffer< u32 > &	buf_reduc_index_map,
		sham::DeviceBuffer< u_morton > &	buf_morton,
		sham::DeviceBuffer< u_morton > &	buf_leaf_morton
	)

Remaps a Morton tree on device using a reduction index map.

This function takes a device buffer buf_reduc_index_map containing the reduction index map and a device buffer buf_morton containing the Morton tree. It will remap the Morton tree using the reduction index map and write the result to a device buffer buf_leaf_morton.

The function takes a device scheduler queue as input and will use it to schedule the kernel calls.

Parameters

queue	a device scheduler to schedule the kernel calls
morton_leaf_count	the number of leaf nodes in the reduced tree
buf_reduc_index_map	a device buffer containing the reduction index map
buf_morton	a device buffer containing the Morton tree
buf_leaf_morton	a device buffer that will contain the remapped Morton tree

Definition at line 565 of file reduction_alg.cpp.

Here is the call graph for this function:

sycl_morton_remap_reduction< u32 >()

template<>

void sycl_morton_remap_reduction< u32 >	(	sycl::queue &	queue,
		u32	morton_leaf_count,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_morton,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_leaf_morton
	)

Definition at line 539 of file reduction_alg.cpp.

sycl_morton_remap_reduction< u64 >()

template<>

void sycl_morton_remap_reduction< u64 >	(	sycl::queue &	queue,
		u32	morton_leaf_count,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map,
		std::unique_ptr< sycl::buffer< u64 > > &	buf_morton,
		std::unique_ptr< sycl::buffer< u64 > > &	buf_leaf_morton
	)

Definition at line 552 of file reduction_alg.cpp.

sycl_reduction_iteration() [1/2]

template<class u_morton , class split_int >

void sycl_reduction_iteration	(	sham::DeviceQueue &	queue,
		u32	morton_count,
		sham::DeviceBuffer< u_morton > &	buf_morton,
		sham::DeviceBuffer< split_int > &	buf_split_table_in,
		sham::DeviceBuffer< split_int > &	buf_split_table_out
	)

Definition at line 197 of file reduction_alg.cpp.

sycl_reduction_iteration() [2/2]

template<class u_morton , class kername , class split_int >

void sycl_reduction_iteration	(	sycl::queue &	queue,
		u32	morton_count,
		std::unique_ptr< sycl::buffer< u_morton > > &	buf_morton,
		std::unique_ptr< sycl::buffer< split_int > > &	buf_split_table_in,
		std::unique_ptr< sycl::buffer< split_int > > &	buf_split_table_out
	)

Definition at line 144 of file reduction_alg.cpp.

update_morton_buf() [1/2]

void update_morton_buf	(	sham::DeviceQueue &	queue,
		u32	len,
		u32	val_ins,
		sham::DeviceBuffer< u32 > &	buf_src,
		sham::DeviceBuffer< u32 > &	buf_reduc_index_map
	)

Definition at line 274 of file reduction_alg.cpp.

update_morton_buf() [2/2]

void update_morton_buf	(	sycl::queue &	queue,
		u32	len,
		u32	val_ins,
		sycl::buffer< u32 > &	buf_src,
		std::unique_ptr< sycl::buffer< u32 > > &	buf_reduc_index_map
	)

Definition at line 246 of file reduction_alg.cpp.