boost/compute
2
0
Fork 0
mirror of https://github.com/boostorg/compute.git synced 2026-02-19 14:22:12 +00:00
Code Issues Projects Releases Wiki Activity

Refactor dispatch_sort() function

Browse Source
This commit is contained in:
Kyle Lutz
2014-12-02 21:49:51 -08:00
parent bb19d2a07c
commit faddbf2368
2 changed files with 75 additions and 157 deletions
Download Patch File Download Diff File Expand all files Collapse all files

126
include/boost/compute/algorithm/detail/fixed_sort.hpp
View File

@@ -1,126 +0,0 @@
//---------------------------------------------------------------------------//
// Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
//
// Distributed under the Boost Software License, Version 1.0
// See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt
//
// See http://kylelutz.github.com/compute for more information.
//---------------------------------------------------------------------------//
#ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_FIXED_SORT_HPP
#define BOOST_COMPUTE_ALGORITHM_DETAIL_FIXED_SORT_HPP
#include <boost/compute/kernel.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/type_traits/type_name.hpp>
#include <boost/compute/detail/program_cache.hpp>
namespace boost {
namespace compute {
namespace detail {
// sort two values
template<class T>
inline void sort2(const buffer &buffer, command_queue &queue)
{
const context &context = queue.get_context();
boost::shared_ptr<detail::program_cache> cache =
detail::get_program_cache(context);
std::string cache_key =
std::string("fixed_sort2_") + type_name<T>();
program sort2_program = cache->get(cache_key);
if(!sort2_program.get()){
const char source[] =
"__kernel void sort2(__global T *input)\n"
"{\n"
" const T x = input[0];\n"
" const T y = input[1];\n"
" if(y < x){\n"
" input[0] = y;\n"
" input[1] = x;\n"
" }\n"
"}\n";
sort2_program = program::build_with_source(
source, context, std::string("-DT=") + type_name<T>()
);
cache->insert(cache_key, sort2_program);
}
kernel sort2_kernel = sort2_program.create_kernel("sort2");
sort2_kernel.set_arg(0, buffer);
queue.enqueue_task(sort2_kernel);
}
// sort three values
template<class T>
inline void sort3(const buffer &buffer, command_queue &queue)
{
const context &context = queue.get_context();
boost::shared_ptr<detail::program_cache> cache =
detail::get_program_cache(context);
std::string cache_key =
std::string("fixed_sort3_") + type_name<T>();
program sort3_program = cache->get(cache_key);
if(!sort3_program.get()){
const char source[] =
"__kernel void sort3(__global T *input)\n"
"{\n"
" const T x = input[0];\n"
" const T y = input[1];\n"
" const T z = input[2];\n"
" if(y < x){\n"
" if(z < x){\n"
" if(z < y){\n"
" input[0] = z;\n"
" input[1] = y;\n"
" input[2] = x;\n"
" }\n"
" else {\n"
" input[0] = y;\n"
" input[1] = z;\n"
" input[2] = x;\n"
" }\n"
" }\n"
" else {\n"
" input[0] = y;\n"
" input[1] = x;\n"
" }\n"
" }\n"
" else {\n"
" if(z < x){\n"
" input[0] = z;\n"
" input[1] = x;\n"
" input[2] = y;\n"
" }\n"
" else if(z < y){\n"
" input[1] = z;\n"
" input[2] = y;\n"
" }\n"
" }\n"
"}\n";
sort3_program = program::build_with_source(
source, context, std::string("-DT=") + type_name<T>()
);
cache->insert(cache_key, sort3_program);
}
kernel sort3_kernel = sort3_program.create_kernel("sort3");
sort3_kernel.set_arg(0, buffer);
queue.enqueue_task(sort3_kernel);
}
} // end detail namespace
} // end compute namespace
} // end boost namespace
#endif // BOOST_COMPUTE_ALGORITHM_DETAIL_FIXED_SORT_HPP

106
include/boost/compute/algorithm/sort.hpp
View File

@@ -15,40 +15,34 @@
#include <boost/utility/enable_if.hpp>
#include <boost/compute/buffer.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/detail/fixed_sort.hpp>
#include <boost/compute/algorithm/detail/radix_sort.hpp>
#include <boost/compute/algorithm/detail/insertion_sort.hpp>
#include <boost/compute/algorithm/reverse.hpp>
#include <boost/compute/container/mapped_view.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/iterator/buffer_iterator.hpp>
namespace boost {
namespace compute {
namespace detail {
// sort() for device iterators
template <class Iterator>
inline void dispatch_sort(Iterator first,
Iterator last,
command_queue &queue,
typename boost::enable_if<
is_device_iterator<Iterator>
>::type* = 0)
template<class T>
inline void dispatch_device_sort(buffer_iterator<T> first,
buffer_iterator<T> last,
less<T>,
command_queue &queue,
typename boost::enable_if_c<
is_radix_sortable<T>::value
>::type* = 0)
{
typedef typename std::iterator_traits<Iterator>::value_type T;
size_t count = detail::iterator_range_size(first, last);
if(count < 2){
// nothing to do
return;
}
else if(count == 2){
::boost::compute::detail::sort2<T>(first.get_buffer(), queue);
}
else if(count == 3){
::boost::compute::detail::sort3<T>(first.get_buffer(), queue);
}
else if(count <= 32){
::boost::compute::detail::serial_insertion_sort(first, last, queue);
}
@@ -57,10 +51,64 @@ inline void dispatch_sort(Iterator first,
}
}
// sort() for host iterators
template <class Iterator>
template<class T>
inline void dispatch_device_sort(buffer_iterator<T> first,
buffer_iterator<T> last,
greater<T> compare,
command_queue &queue,
typename boost::enable_if_c<
is_radix_sortable<T>::value
>::type* = 0)
{
size_t count = detail::iterator_range_size(first, last);
if(count < 2){
// nothing to do
return;
}
else if(count <= 32){
::boost::compute::detail::serial_insertion_sort(
first, last, compare, queue
);
}
else {
// radix sort in ascending order
::boost::compute::detail::radix_sort(first, last, queue);
// reverse range to descending order
::boost::compute::reverse(first, last, queue);
}
}
template<class Iterator, class Compare>
inline void dispatch_device_sort(Iterator first,
Iterator last,
Compare compare,
command_queue &queue)
{
::boost::compute::detail::serial_insertion_sort(
first, last, compare, queue
);
}
// sort() for device iterators
template<class Iterator, class Compare>
inline void dispatch_sort(Iterator first,
Iterator last,
Compare compare,
command_queue &queue,
typename boost::enable_if<
is_device_iterator<Iterator>
>::type* = 0)
{
dispatch_device_sort(first, last, compare, queue);
}
// sort() for host iterators
template<class Iterator, class Compare>
inline void dispatch_sort(Iterator first,
Iterator last,
Compare compare,
command_queue &queue,
typename boost::disable_if<
is_device_iterator<Iterator>
@@ -76,7 +124,7 @@ inline void dispatch_sort(Iterator first,
);
// sort mapped buffer
dispatch_sort(view.begin(), view.end(), queue);
dispatch_device_sort(view.begin(), view.end(), compare, queue);
// return results to host
view.map(queue);
@@ -118,15 +166,7 @@ inline void sort(Iterator first,
Compare compare,
command_queue &queue = system::default_queue())
{
size_t count = detail::iterator_range_size(first, last);
if(count < 2){
return;
}
return ::boost::compute::detail::serial_insertion_sort(first,
last,
compare,
queue);
::boost::compute::detail::dispatch_sort(first, last, compare, queue);
}
/// \overload
@@ -135,7 +175,11 @@ inline void sort(Iterator first,
Iterator last,
command_queue &queue = system::default_queue())
{
detail::dispatch_sort(first, last, queue);
typedef typename std::iterator_traits<Iterator>::value_type value_type;
::boost::compute::sort(
first, last, ::boost::compute::less<value_type>(), queue
);
}
} // end compute namespace