bitonic.cuh File Reference

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

Include dependency graph for bitonic.cuh:

This graph shows which files directly or indirectly include this file:

Go to the source code of this file.

Classes
struct	OrderWithCmpSwap< T >
	Wrapped compare and swap-to-correct-order function. More...
struct	OrderWithMinMax< T >
	Min-max-into-ordering function. More...
struct	ComparatorPolicy< T >
	Comparator policy that automatically chooses a good implementation. More...
struct	ComparatorPolicy< float >

Macros
#define	ESOM_CUDA_BITONIC_CUH

Functions
template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>, bool UP = false>
__device__ void	bitonic_merge_step (T *__restrict__ data)
	A single "layer" of the parallel bitonic comparator. More...
template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>, bool UP = true>
__device__ void	bitonic_merge (T *__restrict__ data)
	Parallel bitonic merge. More...
template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>>
__device__ __forceinline__ void	bitonic_sort (T *__restrict__ data)
	Perform multiple bitonic sorts by all active threads. More...
template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>>
__device__ __forceinline__ void	bitonic_topk_update_opt (T *__restrict__ topK, T *__restrict__ newData)
	Perform one update step of bitonic topk algorithm. More...

Macro Definition Documentation

ESOM_CUDA_BITONIC_CUH

#define ESOM_CUDA_BITONIC_CUH

Definition at line 30 of file bitonic.cuh.

Function Documentation

bitonic_merge()

template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>, bool UP = true>

__device__ void bitonic_merge

(

T *__restrict__

data

)

Parallel bitonic merge.

This runs several layers of the parallel bitonic comparators.

Definition at line 121 of file bitonic.cuh.

Here is the call graph for this function:

Here is the caller graph for this function:

bitonic_merge_step()

template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>, bool UP = false>

__device__ void bitonic_merge_step

(

T *__restrict__

data

)

A single "layer" of the parallel bitonic comparator.

Definition at line 83 of file bitonic.cuh.

bitonic_sort()

template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>>

__device__ __forceinline__ void bitonic_sort

(

T *__restrict__

data

)

Perform multiple bitonic sorts by all active threads.

Template Parameters

T	data item type
BLOCK_SIZE	Number of threads that work cooperatively on a block of items that has exactly 2x BLOCK_SIZE items. Size of thread block must be multiple of BLOCK_SIZE (multiple data blocks may be sorted by one thread block).
CMP	comparator policy class that implements compare and swap on type T

Parameters

data	pointer to an array of T to be sorted in-place, of size at least two times the total number of threads

Definition at line 152 of file bitonic.cuh.

Here is the call graph for this function:

Here is the caller graph for this function:

bitonic_topk_update_opt()

template<typename T , int BLOCK_SIZE, class CMP = ComparatorPolicy<T>>

__device__ __forceinline__ void bitonic_topk_update_opt	(	T *__restrict__	topK,
		T *__restrict__	newData
	)

Perform one update step of bitonic topk algorithm.

The algorith takes two inputs: current topk sub-result and new data (e.g., newly computed distances). It sorts inputs and runs a bitonic merge on these, producing a bitonic sequence with the lower half of the data in topK, and (as a side product) a bitonic sequence with the upper half of the data in newData.

Note that hat way the topk part is updated, but not entirely sorted (so we save some time).

Template Parameters

T	item data type
BLOCK_SIZE	Number of threads that work cooperatively on a block of items. Both topk result and newData of one block have 2x BLOCK_SIZE items. Size of thread block must be multiple of BLOCK_SIZE (multiple data blocks may be sorted by one thread block).
CMP	comparator policy, such as ComparatorPolicy.

Parameters

topK	array of T's with intermediate top results, containing at least 2x(number of threads) items
newData	array of new results to be merged into topK, of the same size

Definition at line 198 of file bitonic.cuh.