namespace to store algorithms implemented by shamalgs More...

Classes
class	DigitBinner
class	SortByKeyRadixOnesweep
struct	OddEvenOrderingPrimitive
	Device side primitives of the odd-even merge network. More...
struct	OrderingPrimitive
struct	OrderingPrimitiveXorSwap

Functions
template<class Tkey, class Tval>
void	sort_by_key_batcher_odd_even (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< Tkey > &buf_key, sham::DeviceBuffer< Tval > &buf_values, u32 len)
	Sort key-value pairs of any length using a Batcher odd-even merge network.
template<class Tkey, class Tval>
void	sort_by_key_batcher_odd_even_host_reference (std::vector< Tkey > &keys, std::vector< Tval > &values)
	Host reference of sort_by_key_batcher_odd_even.
template<class Tkey, class Tval>
void	sort_by_key_bitonic_legacy (sycl::queue &q, sycl::buffer< Tkey > &buf_key, sycl::buffer< Tval > &buf_values, u32 len)
template<class Tkey, class Tval, u32 MaxStencilSize>
void	sort_by_key_bitonic_updated (sycl::queue &q, sycl::buffer< Tkey > &buf_key, sycl::buffer< Tval > &buf_values, u32 len)
template<class Tkey, class Tval, u32 MaxStencilSize>
void	sort_by_key_bitonic_updated_xor_swap (sycl::queue &q, sycl::buffer< Tkey > &buf_key, sycl::buffer< Tval > &buf_values, u32 len)
template<class Tkey, class Tval>
void	sort_by_key_bitonic_fallback (sycl::queue &q, sycl::buffer< Tkey > &buf_key, sycl::buffer< Tval > &buf_values, u32 len)
template<class Tkey, class Tval, u32 MaxStencilSize>
void	sort_by_key_bitonic_updated_usm (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< Tkey > &buf_key, sham::DeviceBuffer< Tval > &buf_values, u32 len)
template<class Tkey, class Tval, u32 group_size, u32 digit_len>
void	sort_by_key_radix_onesweep (sycl::queue &q, sycl::buffer< Tkey > &buf_key, sycl::buffer< Tval > &buf_values, u32 len)
template void	sort_by_key_batcher_odd_even< u32, u32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< u32 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_batcher_odd_even< u64, u32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< u64 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_batcher_odd_even< f32, f32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f32 > &buf_key, sham::DeviceBuffer< f32 > &buf_values, u32 len)
template void	sort_by_key_batcher_odd_even< f64, f64 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f64 > &buf_key, sham::DeviceBuffer< f64 > &buf_values, u32 len)
template void	sort_by_key_batcher_odd_even< f32, u32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f32 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_batcher_odd_even< f64, u32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f64 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_batcher_odd_even_host_reference< u32, u32 > (std::vector< u32 > &keys, std::vector< u32 > &values)
template void	sort_by_key_batcher_odd_even_host_reference< u64, u32 > (std::vector< u64 > &keys, std::vector< u32 > &values)
template void	sort_by_key_batcher_odd_even_host_reference< f32, f32 > (std::vector< f32 > &keys, std::vector< f32 > &values)
template void	sort_by_key_batcher_odd_even_host_reference< f32, u32 > (std::vector< f32 > &keys, std::vector< u32 > &values)
template void	sort_by_key_batcher_odd_even_host_reference< f64, u32 > (std::vector< f64 > &keys, std::vector< u32 > &values)
template void	sort_by_key_batcher_odd_even_host_reference< f64, f64 > (std::vector< f64 > &keys, std::vector< f64 > &values)
template void	sort_by_key_bitonic_legacy (sycl::queue &q, sycl::buffer< u32 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_legacy (sycl::queue &q, sycl::buffer< u64 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated< u32, u32, 16 > (sycl::queue &q, sycl::buffer< u32 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated< u64, u32, 16 > (sycl::queue &q, sycl::buffer< u64 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated< u32, u32, 8 > (sycl::queue &q, sycl::buffer< u32 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated< u64, u32, 8 > (sycl::queue &q, sycl::buffer< u64 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated< u32, u32, 32 > (sycl::queue &q, sycl::buffer< u32 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated< u64, u32, 32 > (sycl::queue &q, sycl::buffer< u64 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< u32, u32, 16 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< u32 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< u32, u32, 32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< u32 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< u64, u32, 16 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< u64 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< u64, u32, 32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< u64 > &buf_key, sham::DeviceBuffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< f32, f32, 16 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f32 > &buf_key, sham::DeviceBuffer< f32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< f32, f32, 32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f32 > &buf_key, sham::DeviceBuffer< f32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< f64, f64, 16 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f64 > &buf_key, sham::DeviceBuffer< f64 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_usm< f64, f64, 32 > (const sham::DeviceScheduler_ptr &sched, sham::DeviceBuffer< f64 > &buf_key, sham::DeviceBuffer< f64 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_xor_swap< u32, u32, 16 > (sycl::queue &q, sycl::buffer< u32 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_xor_swap< u64, u32, 16 > (sycl::queue &q, sycl::buffer< u64 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_xor_swap< u32, u32, 8 > (sycl::queue &q, sycl::buffer< u32 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_xor_swap< u64, u32, 8 > (sycl::queue &q, sycl::buffer< u64 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_xor_swap< u32, u32, 32 > (sycl::queue &q, sycl::buffer< u32 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)
template void	sort_by_key_bitonic_updated_xor_swap< u64, u32, 32 > (sycl::queue &q, sycl::buffer< u64 > &buf_key, sycl::buffer< u32 > &buf_values, u32 len)

Detailed Description

namespace to store algorithms implemented by shamalgs

Function Documentation

◆ sort_by_key_batcher_odd_even()

template<class Tkey, class Tval>

void shamalgs::algorithm::details::sort_by_key_batcher_odd_even	(	const sham::DeviceScheduler_ptr &	sched,
		sham::DeviceBuffer< Tkey > &	buf_key,
		sham::DeviceBuffer< Tval > &	buf_values,
		u32	len )

Sort key-value pairs of any length using a Batcher odd-even merge network.

Both buffers are modified in place: the keys are sorted in ascending order and the values follow the same permutation. Unlike the bitonic implementations, len is unconstrained, and no padding buffer is allocated.

Template Parameters

Tkey	Key type, must be comparable (supports operator<)
Tval	Value type, must be copyable

Parameters

sched	Device scheduler used for the kernel launches
buf_key	Device buffer holding the keys to sort by
buf_values	Device buffer holding the values to reorder
len	Length of both buffers, any value including 0 and 1

Note: The sort is not stable, equal keys may be reordered; One kernel is launched per stage of the network, ceil(log2(len))*(ceil(log2(len))+1)/2 in total, each with ceil(len/2) threads

Definition at line 99 of file batcherOddEvenSort.cpp.

Here is the call graph for this function:

◆ sort_by_key_batcher_odd_even_host_reference()

template<class Tkey, class Tval>

void shamalgs::algorithm::details::sort_by_key_batcher_odd_even_host_reference	(	std::vector< Tkey > &	keys,
		std::vector< Tval > &	values )

Host reference of sort_by_key_batcher_odd_even.

A literal transcription of the network quoted at the top of this file, written as four plain loops with no index algebra. It is the oracle the device implementation is tested against: both run the very same comparators in the very same order, so their outputs must match element for element even though the sort is unstable.

Template Parameters

Tkey	Key type, must be comparable (supports operator<)
Tval	Value type, must be swappable

Parameters

keys	Keys to sort by, sorted in place
values	Values to reorder, permuted in place alongside the keys

Definition at line 140 of file batcherOddEvenSort.cpp.

Here is the call graph for this function:

◆ sort_by_key_bitonic_fallback()

template<class Tkey, class Tval>

void shamalgs::algorithm::details::sort_by_key_bitonic_fallback	(	sycl::queue &	q,
		sycl::buffer< Tkey > &	buf_key,
		sycl::buffer< Tval > &	buf_values,
		u32	len )

inline

Definition at line 37 of file bitonicSort.hpp.

◆ sort_by_key_bitonic_legacy()

template<class Tkey, class Tval>

void shamalgs::algorithm::details::sort_by_key_bitonic_legacy	(	sycl::queue &	q,
		sycl::buffer< Tkey > &	buf_key,
		sycl::buffer< Tval > &	buf_values,
		u32	len )

Definition at line 102 of file bitonicSort_legacy.cpp.

◆ sort_by_key_bitonic_updated()

template<class Tkey, class Tval, u32 MaxStencilSize>

void shamalgs::algorithm::details::sort_by_key_bitonic_updated	(	sycl::queue &	q,
		sycl::buffer< Tkey > &	buf_key,
		sycl::buffer< Tval > &	buf_values,
		u32	len )

Definition at line 284 of file bitonicSort_updated.cpp.

◆ sort_by_key_bitonic_updated_usm()

template<class Tkey, class Tval, u32 MaxStencilSize>

void shamalgs::algorithm::details::sort_by_key_bitonic_updated_usm	(	const sham::DeviceScheduler_ptr &	sched,
		sham::DeviceBuffer< Tkey > &	buf_key,
		sham::DeviceBuffer< Tval > &	buf_values,
		u32	len )

Definition at line 285 of file bitonicSort_updated_usm.cpp.

◆ sort_by_key_bitonic_updated_xor_swap()

template<class Tkey, class Tval, u32 MaxStencilSize>

void shamalgs::algorithm::details::sort_by_key_bitonic_updated_xor_swap	(	sycl::queue &	q,
		sycl::buffer< Tkey > &	buf_key,
		sycl::buffer< Tval > &	buf_values,
		u32	len )

Definition at line 263 of file bitonicSort_updated_xor_swap.cpp.

◆ sort_by_key_radix_onesweep()

template<class Tkey, class Tval, u32 group_size, u32 digit_len>

void shamalgs::algorithm::details::sort_by_key_radix_onesweep	(	sycl::queue &	q,
		sycl::buffer< Tkey > &	buf_key,
		sycl::buffer< Tval > &	buf_values,
		u32	len )

Definition at line 54 of file radixSortOnesweep.hpp.

Classes

Functions

Detailed Description

Function Documentation

◆ sort_by_key_batcher_odd_even()

◆ sort_by_key_batcher_odd_even_host_reference()

◆ sort_by_key_bitonic_fallback()

◆ sort_by_key_bitonic_legacy()

◆ sort_by_key_bitonic_updated()

◆ sort_by_key_bitonic_updated_usm()

◆ sort_by_key_bitonic_updated_xor_swap()

◆ sort_by_key_radix_onesweep()