Priority Queue

[andrewbriand]

Adds a GPU-accelerated priority queue

Allows for multiple concurrent insertions as well as multiple concurrent
deletions.

The implementation of the priority queue is based on https://arxiv.org/pdf/1906.06504.pdf.

The queue supports two operations:
push: Add elements into the queue
pop: Remove the element(s) with the lowest (when Max == false) or highest
(when Max == true) keys

The priority queue supports bulk host-side operations and more fine-grained
device-side operations.

The host-side bulk operations push and pop allow an arbitrary number of
elements to be pushed to or popped from the queue.

The device-side operations allow a cooperative group to push or pop
some number of elements less than or equal to node_size. These device side
operations are invoked with a trivially-copyable device view,
device_mutable_view which can be obtained with the host function
get_mutable_device_view and passed to the device.

Current limitations:

• Only supports trivially comparable key types
• Does not support insertion and deletion at the same time
• Capacity is fixed and the queue does not automatically resize
• Deletion from the queue is much slower than insertion into the queue due to congestion at the underlying heap's root node

TODO: Port tests to Catch2 and benchmarks to google benchmark

[GPUtester]

Can one of the admins verify this patch?

[jrhemstad]

Is there a reason for this to be hardcoded for key-value pairs? Can't it be for any trivially copyable type T? e.g., with std::priority_queue I could have a std::priority_queue<int> or a std::priority_queue<std::pair<int,int>>.

[jrhemstad]

Update docs now that this has been updated.

[jrhemstad]

I reviewed the top level header at this point and gave some thoughts/questions on how to make this a little more generic.

[PointKernel]

ok to test

[andrewbriand]

@PointKernel Thanks for your comments! I believe that I have addressed or responded to them all. Please let me know what you think and what other comments you might have.

[PointKernel]

Another round of review.

Thanks @andrewbriand for your effort and persistence made to this PR! We are almost there.

[PointKernel]

Suggested change

	public:
	public:
	using value_type = T;

[andrewbriand]

Should I also replace references to T with value_type in device_mutable_view?

[PointKernel]

That will be great!

[PointKernel]

Suggested change

	detail::push_kernel<<<num_blocks, block_size, get_shmem_size(block_size), stream>>>(
	first,
	last - first,
	d_heap_,
	d_size_,
	node_size_,
	d_locks_,
	d_p_buffer_size_,
	lowest_level_start_,
	compare_);
	auto view = get_device_mutable_view();
	detail::push_kernel<<<num_blocks, block_size, get_shmem_size(block_size), stream>>>(
	first, num_elements, view);

This is a great example showing the power of "view". Accordingly, the push_kernel would look like:

template <typename OutputIt, typename viewT>
__global__ void push_kernel(OutputIt elements,
                            std::size_t const num_elements,
                            viewT view)
{
  using T = typename viewT::value_type;
  ...
}

If you want, push_n_kernel instead of push_kernel would be a more descriptive name in this case.

[PointKernel]

push_single_node, push_partial_node, and related utilities should be member functions of device_mutable_view.

[PointKernel]

The same as pop_single_node and pop_partial_node

[PointKernel]

This seems never used

[PointKernel]

OutputIt instead of InputIt

[PointKernel]

Suggested change

	template <typename OutputIt, typename T, typename Compare>
	template <typename InputIt, typename viewT>

[PointKernel]

Suggested change

	T* heap,
	int* size,
	std::size_t node_size,
	int* locks,
	std::size_t* p_buffer_size,
	int lowest_level_start,
	Compare compare)
	viewT view)

The kernel implementation can also be simplified with view.

[PointKernel]

@andrewbriand Can you please also merge with the latest dev branch and fix build warnings (if there is any)?