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

Merge remote-tracking branch 'compute/develop'

Browse Source
This commit is contained in:
Anees Haider
2015-04-06 09:33:52 +05:00
parent 31342d7e53 72c75f189d
commit b565232b8e
41 changed files with 1424 additions and 493 deletions
Download Patch File Download Diff File Expand all files Collapse all files

90
include/boost/compute/algorithm/copy_if.hpp
View File

@@ -11,87 +11,13 @@
#ifndef BOOST_COMPUTE_ALGORITHM_COPY_IF_HPP
#define BOOST_COMPUTE_ALGORITHM_COPY_IF_HPP
#include <boost/compute/cl.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/count.hpp>
#include <boost/compute/algorithm/count_if.hpp>
#include <boost/compute/algorithm/exclusive_scan.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/iterator/discard_iterator.hpp>
#include <boost/compute/algorithm/transform_if.hpp>
#include <boost/compute/functional/identity.hpp>
namespace boost {
namespace compute {
namespace detail {
template<class InputIterator, class OutputIterator, class Predicate>
inline OutputIterator copy_if_impl(InputIterator first,
InputIterator last,
OutputIterator result,
Predicate predicate,
bool copyIndex,
command_queue &queue)
{
typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
size_t count = detail::iterator_range_size(first, last);
if(count == 0){
return result;
}
const context &context = queue.get_context();
// storage for destination indices
::boost::compute::vector<cl_uint> indices(count, context);
// write counts
::boost::compute::detail::meta_kernel k1("copy_if_write_counts");
k1 << indices.begin()[k1.get_global_id(0)] << " = "
<< predicate(first[k1.get_global_id(0)]) << " ? 1 : 0;\n";
k1.exec_1d(queue, 0, count);
// count number of elements to be copied
size_t copied_element_count =
::boost::compute::count(indices.begin(), indices.end(), 1, queue);
// scan indices
::boost::compute::exclusive_scan(indices.begin(),
indices.end(),
indices.begin(),
queue);
// copy values
::boost::compute::detail::meta_kernel k2("copy_if_do_copy");
k2 << "if(" << predicate(first[k2.get_global_id(0)]) << ")" <<
" " << result[indices.begin()[k2.get_global_id(0)]] << "=";
if(copyIndex){
k2 << k2.get_global_id(0) << ";\n";
}
else {
k2 << first[k2.get_global_id(0)] << ";\n";
}
k2.exec_1d(queue, 0, count);
return result + static_cast<difference_type>(copied_element_count);
}
template<class InputIterator, class Predicate>
inline discard_iterator copy_if_impl(InputIterator first,
InputIterator last,
discard_iterator result,
Predicate predicate,
bool copyIndex,
command_queue &queue)
{
(void) copyIndex;
return result + count_if(first, last, predicate, queue);
}
// like the copy_if() algorithm but writes the indices of the values for which
// predicate returns true.
template<class InputIterator, class OutputIterator, class Predicate>
@@ -101,7 +27,11 @@ inline OutputIterator copy_index_if(InputIterator first,
Predicate predicate,
command_queue &queue = system::default_queue())
{
return detail::copy_if_impl(first, last, result, predicate, true, queue);
typedef typename std::iterator_traits<InputIterator>::value_type T;
return detail::transform_if_impl(
first, last, result, identity<T>(), predicate, true, queue
);
}
} // end detail namespace
@@ -115,7 +45,11 @@ inline OutputIterator copy_if(InputIterator first,
Predicate predicate,
command_queue &queue = system::default_queue())
{
return detail::copy_if_impl(first, last, result, predicate, false, queue);
typedef typename std::iterator_traits<InputIterator>::value_type T;
return ::boost::compute::transform_if(
first, last, result, identity<T>(), predicate, queue
);
}
} // end compute namespace

23
include/boost/compute/algorithm/detail/binary_find.hpp
View File

@@ -15,6 +15,7 @@
#include <boost/compute/algorithm/find_if.hpp>
#include <boost/compute/algorithm/transform.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/parameter_cache.hpp>
namespace boost {
namespace compute {
@@ -28,11 +29,9 @@ namespace detail{
class binary_find_kernel : public meta_kernel
{
public:
size_t threads;
binary_find_kernel() : meta_kernel("binary_find")
binary_find_kernel(size_t threads) : meta_kernel("binary_find")
{
threads = 128;
m_threads = threads;
}
template<class InputIterator, class UnaryPredicate>
@@ -41,7 +40,7 @@ public:
UnaryPredicate predicate)
{
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
int block = (iterator_range_size(first, last)-1)/(threads-1);
int block = (iterator_range_size(first, last)-1)/(m_threads-1);
m_index_arg = add_arg<uint_ *>(memory_object::global_memory, "index");
@@ -60,10 +59,11 @@ public:
{
set_arg(m_index_arg, index.get_buffer());
return exec_1d(queue, 0, threads);
return exec_1d(queue, 0, m_threads);
}
private:
size_t m_threads;
size_t m_index_arg;
};
@@ -84,8 +84,15 @@ inline InputIterator binary_find(InputIterator first,
UnaryPredicate predicate,
command_queue &queue = system::default_queue())
{
const device &device = queue.get_device();
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
const std::string cache_key = "__boost_binary_find";
size_t find_if_limit = 128;
size_t threads = 128;
size_t threads = parameters->get(cache_key, "tpb", 128);
size_t count = iterator_range_size(first, last);
while(count > find_if_limit) {
@@ -93,7 +100,7 @@ inline InputIterator binary_find(InputIterator first,
scalar<uint_> index(queue.get_context());
index.write(static_cast<uint_>(count), queue);
binary_find_kernel kernel;
binary_find_kernel kernel(threads);
kernel.set_range(first, last, predicate);
kernel.exec(queue, index);

24
include/boost/compute/algorithm/detail/copy_on_device.hpp
View File

@@ -20,6 +20,7 @@
#include <boost/compute/memory/svm_ptr.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/parameter_cache.hpp>
#include <boost/compute/detail/work_size.hpp>
namespace boost {
@@ -42,12 +43,21 @@ template<class InputIterator, class OutputIterator>
class copy_kernel : public meta_kernel
{
public:
copy_kernel()
copy_kernel(const device &device)
: meta_kernel("copy")
{
m_count = 0;
m_vpt = 4;
m_tpb = 128;
typedef typename std::iterator_traits<InputIterator>::value_type input_type;
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
std::string cache_key =
"__boost_copy_kernel_" + boost::lexical_cast<std::string>(sizeof(input_type));
m_vpt = parameters->get(cache_key, "vpt", 4);
m_tpb = parameters->get(cache_key, "tpb", 128);
}
void set_range(InputIterator first,
@@ -97,7 +107,9 @@ inline OutputIterator copy_on_device(InputIterator first,
OutputIterator result,
command_queue &queue)
{
copy_kernel<InputIterator, OutputIterator> kernel;
const device &device = queue.get_device();
copy_kernel<InputIterator, OutputIterator> kernel(device);
kernel.set_range(first, last, result);
kernel.exec(queue);
@@ -122,7 +134,9 @@ inline future<OutputIterator> copy_on_device_async(InputIterator first,
OutputIterator result,
command_queue &queue)
{
copy_kernel<InputIterator, OutputIterator> kernel;
const device &device = queue.get_device();
copy_kernel<InputIterator, OutputIterator> kernel(device);
kernel.set_range(first, last, result);
event event_ = kernel.exec(queue);

2
include/boost/compute/algorithm/detail/inplace_reduce.hpp
View File

@@ -13,6 +13,8 @@
#include <iterator>
#include <boost/utility/result_of.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>

43
include/boost/compute/algorithm/detail/radix_sort.hpp
View File

@@ -23,9 +23,10 @@
#include <boost/compute/algorithm/exclusive_scan.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/detail/parameter_cache.hpp>
#include <boost/compute/type_traits/type_name.hpp>
#include <boost/compute/type_traits/is_fundamental.hpp>
#include <boost/compute/type_traits/is_vector_type.hpp>
#include <boost/compute/type_traits/type_name.hpp>
#include <boost/compute/utility/program_cache.hpp>
namespace boost {
@@ -232,19 +233,9 @@ inline void radix_sort_impl(const buffer_iterator<T> first,
typedef T value_type;
typedef typename radix_sort_value_type<sizeof(T)>::type sort_type;
const device &device = queue.get_device();
const context &context = queue.get_context();
size_t count = detail::iterator_range_size(first, last);
// sort parameters
const uint_ k = 4;
const uint_ k2 = 1 << k;
const uint_ block_size = 128;
uint_ block_count = static_cast<uint_>(count / block_size);
if(block_count * block_size != count){
block_count++;
}
// if we have a valid values iterator then we are doing a
// sort by key and have to set up the values buffer
@@ -258,6 +249,17 @@ inline void radix_sort_impl(const buffer_iterator<T> first,
cache_key += std::string("_with_") + type_name<T2>();
}
boost::shared_ptr<program_cache> cache =
program_cache::get_global_cache(context);
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
// sort parameters
const uint_ k = parameters->get(cache_key, "k", 4);
const uint_ k2 = 1 << k;
const uint_ block_size = parameters->get(cache_key, "tpb", 128);
// sort program compiler options
std::stringstream options;
options << "-DK_BITS=" << k;
options << " -DT=" << type_name<sort_type>();
@@ -277,17 +279,22 @@ inline void radix_sort_impl(const buffer_iterator<T> first,
options << enable_double<T2>();
}
// load (or create) radix sort program
boost::shared_ptr<program_cache> cache =
program_cache::get_global_cache(context);
program radix_sort_program =
cache->get_or_build(cache_key, options.str(), radix_sort_source, context);
// load radix sort program
program radix_sort_program = cache->get_or_build(
cache_key, options.str(), radix_sort_source, context
);
kernel count_kernel(radix_sort_program, "count");
kernel scan_kernel(radix_sort_program, "scan");
kernel scatter_kernel(radix_sort_program, "scatter");
size_t count = detail::iterator_range_size(first, last);
uint_ block_count = static_cast<uint_>(count / block_size);
if(block_count * block_size != count){
block_count++;
}
// setup temporary buffers
vector<value_type> output(count, context);
vector<T2> values_output(sort_by_key ? count : 0, context);

37
include/boost/compute/algorithm/detail/reduce_on_gpu.hpp
View File

@@ -13,15 +13,15 @@
#include <iterator>
#include <boost/compute/utility/source.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/vendor.hpp>
#include <boost/compute/detail/parameter_cache.hpp>
#include <boost/compute/detail/work_size.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/type_traits/result_of.hpp>
#include <boost/compute/type_traits/type_name.hpp>
#include <boost/compute/utility/program_cache.hpp>
#include <boost/compute/utility/source.hpp>
namespace boost {
namespace compute {
@@ -99,7 +99,7 @@ inline void initial_reduce(InputIterator first,
(void) reduce_kernel;
typedef typename std::iterator_traits<InputIterator>::value_type Arg;
typedef typename ::boost::compute::result_of<Function(Arg, Arg)>::type T;
typedef typename boost::tr1_result_of<Function(Arg, Arg)>::type T;
size_t count = std::distance(first, last);
detail::meta_kernel k("initial_reduce");
@@ -174,6 +174,7 @@ inline void reduce_on_gpu(InputIterator first,
command_queue &queue)
{
const device &device = queue.get_device();
const context &context = queue.get_context();
detail::meta_kernel k("reduce");
k.add_arg<const T*>(memory_object::global_memory, "input");
@@ -210,28 +211,34 @@ inline void reduce_on_gpu(InputIterator first,
" output[output_offset + get_group_id(0)] = scratch[0];\n" <<
"}\n";
uint_ vpt = 8;
uint_ tpb = 128;
std::string cache_key = std::string("__boost_reduce_on_gpu_") + type_name<T>();
size_t count = std::distance(first, last);
// load parameters
boost::shared_ptr<parameter_cache> parameters =
detail::parameter_cache::get_global_cache(device);
const context &context = queue.get_context();
uint_ vpt = parameters->get(cache_key, "vpt", 8);
uint_ tpb = parameters->get(cache_key, "tpb", 128);
// load (or create) reduce program
// reduce program compiler flags
std::stringstream options;
options << "-DT=" << type_name<T>()
<< " -DVPT=" << vpt
<< " -DTPB=" << tpb;
// load program
boost::shared_ptr<program_cache> cache =
program_cache::get_global_cache(context);
std::string cache_key = std::string("__boost_reduce_on_gpu_") + type_name<T>();
std::stringstream options;
options << "-DT=" << type_name<T>() << " -DVPT=" << vpt << " -DTPB=" << tpb;
program reduce_program =
cache->get_or_build(cache_key, options.str(), k.source(), context);
program reduce_program = cache->get_or_build(
cache_key, options.str(), k.source(), context
);
// create reduce kernel
kernel reduce_kernel(reduce_program, "reduce");
size_t count = std::distance(first, last);
// first pass, reduce from input to ping
buffer ping(context, std::ceil(float(count) / vpt / tpb) * sizeof(T));
initial_reduce(first, last, ping, function, reduce_kernel, vpt, tpb, queue);

35
include/boost/compute/algorithm/reverse_copy.hpp
View File

@@ -20,6 +20,33 @@
namespace boost {
namespace compute {
namespace detail {
template<class Iterator, class OutputIterator>
struct reverse_copy_kernel : public meta_kernel
{
reverse_copy_kernel(Iterator first, Iterator last, OutputIterator result)
: meta_kernel("reverse_copy")
{
// store size of the range
m_size = detail::iterator_range_size(first, last);
add_set_arg<const cl_uint>("size", static_cast<const cl_uint>(m_size));
*this <<
decl<cl_uint>("i") << " = get_global_id(0);\n" <<
decl<cl_uint>("j") << " = size - get_global_id(0) - 1;\n" <<
result[var<cl_uint>("j")] << "=" << first[var<cl_uint>("i")] << ";\n";
}
void exec(command_queue &queue)
{
exec_1d(queue, 0, m_size);
}
size_t m_size;
};
} // end detail namespace
/// Copies the elements in the range [\p first, \p last) in reversed
/// order to the range beginning at \p result.
@@ -36,11 +63,11 @@ reverse_copy(InputIterator first,
difference_type count = std::distance(first, last);
// copy data to result
::boost::compute::copy(first, last, result, queue);
detail::reverse_copy_kernel<InputIterator, OutputIterator>
kernel(first, last, result);
// reverse result
::boost::compute::reverse(result, result + count, queue);
// run kernel
kernel.exec(queue);
// return iterator to the end of result
return result + count;

117
include/boost/compute/algorithm/transform_if.hpp Normal file
View File

@@ -0,0 +1,117 @@
//---------------------------------------------------------------------------//
// Copyright (c) 2013-2015 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_TRANSFORM_IF_HPP
#define BOOST_COMPUTE_ALGORITHM_TRANSFORM_IF_HPP
#include <boost/compute/cl.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/command_queue.hpp>
#include <boost/compute/algorithm/count.hpp>
#include <boost/compute/algorithm/count_if.hpp>
#include <boost/compute/algorithm/exclusive_scan.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/iterator/discard_iterator.hpp>
namespace boost {
namespace compute {
namespace detail {
template<class InputIterator, class OutputIterator, class UnaryFunction, class Predicate>
inline OutputIterator transform_if_impl(InputIterator first,
InputIterator last,
OutputIterator result,
UnaryFunction function,
Predicate predicate,
bool copyIndex,
command_queue &queue)
{
typedef typename std::iterator_traits<OutputIterator>::difference_type difference_type;
size_t count = detail::iterator_range_size(first, last);
if(count == 0){
return result;
}
const context &context = queue.get_context();
// storage for destination indices
::boost::compute::vector<cl_uint> indices(count, context);
// write counts
::boost::compute::detail::meta_kernel k1("transform_if_write_counts");
k1 << indices.begin()[k1.get_global_id(0)] << " = "
<< predicate(first[k1.get_global_id(0)]) << " ? 1 : 0;\n";
k1.exec_1d(queue, 0, count);
// count number of elements to be copied
size_t copied_element_count =
::boost::compute::count(indices.begin(), indices.end(), 1, queue);
// scan indices
::boost::compute::exclusive_scan(
indices.begin(), indices.end(), indices.begin(), queue
);
// copy values
::boost::compute::detail::meta_kernel k2("transform_if_do_copy");
k2 << "if(" << predicate(first[k2.get_global_id(0)]) << ")" <<
" " << result[indices.begin()[k2.get_global_id(0)]] << "=";
if(copyIndex){
k2 << k2.get_global_id(0) << ";\n";
}
else {
k2 << function(first[k2.get_global_id(0)]) << ";\n";
}
k2.exec_1d(queue, 0, count);
return result + static_cast<difference_type>(copied_element_count);
}
template<class InputIterator, class UnaryFunction, class Predicate>
inline discard_iterator transform_if_impl(InputIterator first,
InputIterator last,
discard_iterator result,
UnaryFunction function,
Predicate predicate,
bool copyIndex,
command_queue &queue)
{
(void) function;
(void) copyIndex;
return result + count_if(first, last, predicate, queue);
}
} // end detail namespace
/// Copies each element in the range [\p first, \p last) for which
/// \p predicate returns \c true to the range beginning at \p result.
template<class InputIterator, class OutputIterator, class UnaryFunction, class Predicate>
inline OutputIterator transform_if(InputIterator first,
InputIterator last,
OutputIterator result,
UnaryFunction function,
Predicate predicate,
command_queue &queue = system::default_queue())
{
return detail::transform_if_impl(
first, last, result, function, predicate, false, queue
);
}
} // end compute namespace
} // end boost namespace
#endif // BOOST_COMPUTE_ALGORITHM_TRANSFORM_IF_HPP

3
include/boost/compute/detail/get_object_info.hpp
View File

@@ -14,6 +14,9 @@
#include <string>
#include <vector>
#include <boost/preprocessor/seq/for_each.hpp>
#include <boost/preprocessor/tuple/elem.hpp>
#include <boost/throw_exception.hpp>
#include <boost/compute/cl.hpp>

215
include/boost/compute/detail/parameter_cache.hpp Normal file
View File

@@ -0,0 +1,215 @@
//---------------------------------------------------------------------------//
// Copyright (c) 2013-2015 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_DETAIL_PARAMETER_CACHE_HPP
#define BOOST_COMPUTE_DETAIL_PARAMETER_CACHE_HPP
#include <algorithm>
#include <string>
#include <boost/shared_ptr.hpp>
#include <boost/make_shared.hpp>
#include <boost/noncopyable.hpp>
#include <boost/compute/config.hpp>
#include <boost/compute/device.hpp>
#include <boost/compute/detail/global_static.hpp>
#include <boost/compute/version.hpp>
#ifdef BOOST_COMPUTE_USE_OFFLINE_CACHE
#include <boost/algorithm/string/trim.hpp>
#include <boost/compute/detail/path.hpp>
#include <boost/property_tree/ptree.hpp>
#include <boost/property_tree/json_parser.hpp>
#endif // BOOST_COMPUTE_USE_OFFLINE_CACHE
namespace boost {
namespace compute {
namespace detail {
class parameter_cache : boost::noncopyable
{
public:
parameter_cache(const device &device)
: m_dirty(false),
m_device_name(device.name())
{
#ifdef BOOST_COMPUTE_USE_OFFLINE_CACHE
// get offline cache file name (e.g. /home/user/.boost_compute/tune/device.json)
m_file_name = make_file_name();
// load parameters from offline cache file (if it exists)
if(boost::filesystem::exists(m_file_name)){
read_from_disk();
}
#endif // BOOST_COMPUTE_USE_OFFLINE_CACHE
}
~parameter_cache()
{
#ifdef BOOST_COMPUTE_USE_OFFLINE_CACHE
write_to_disk();
#endif // BOOST_COMPUTE_USE_OFFLINE_CACHE
}
void set(const std::string &object, const std::string &parameter, uint_ value)
{
m_cache[std::make_pair(object, parameter)] = value;
// set the dirty flag to true. this will cause the updated parameters
// to be stored to disk.
m_dirty = true;
}
uint_ get(const std::string &object, const std::string &parameter, uint_ default_value)
{
std::map<std::pair<std::string, std::string>, uint_>::iterator
iter = m_cache.find(std::make_pair(object, parameter));
if(iter != m_cache.end()){
return iter->second;
}
else {
return default_value;
}
}
static boost::shared_ptr<parameter_cache> get_global_cache(const device &device)
{
// device name -> parameter cache
typedef std::map<std::string, boost::shared_ptr<parameter_cache> > cache_map;
BOOST_COMPUTE_DETAIL_GLOBAL_STATIC(cache_map, caches, ((std::less<std::string>())));
cache_map::iterator iter = caches.find(device.name());
if(iter == caches.end()){
boost::shared_ptr<parameter_cache> cache =
boost::make_shared<parameter_cache>(device);
caches.insert(iter, std::make_pair(device.name(), cache));
return cache;
}
else {
return iter->second;
}
}
private:
#ifdef BOOST_COMPUTE_USE_OFFLINE_CACHE
// returns a string containing a cannoical device name
static std::string cannonical_device_name(std::string name)
{
boost::algorithm::trim(name);
std::replace(name.begin(), name.end(), ' ', '_');
std::replace(name.begin(), name.end(), '(', '_');
std::replace(name.begin(), name.end(), ')', '_');
return name;
}
// returns the boost.compute version string
static std::string version_string()
{
char buf[32];
std::snprintf(buf, sizeof(buf), "%d.%d.%d", BOOST_COMPUTE_VERSION_MAJOR,
BOOST_COMPUTE_VERSION_MINOR,
BOOST_COMPUTE_VERSION_PATCH);
return buf;
}
// returns the file path for the cached parameters
std::string make_file_name() const
{
return detail::parameter_cache_path(true) + cannonical_device_name(m_device_name) + ".json";
}
// store current parameters to disk
void write_to_disk()
{
BOOST_ASSERT(!m_file_name.empty());
if(m_dirty){
// save current parameters to disk
boost::property_tree::ptree pt;
pt.put("header.device", m_device_name);
pt.put("header.version", version_string());
typedef std::map<std::pair<std::string, std::string>, uint_> map_type;
for(map_type::const_iterator iter = m_cache.begin(); iter != m_cache.end(); ++iter){
const std::pair<std::string, std::string> &key = iter->first;
pt.add(key.first + "." + key.second, iter->second);
}
write_json(m_file_name, pt);
m_dirty = false;
}
}
// load stored parameters from disk
void read_from_disk()
{
BOOST_ASSERT(!m_file_name.empty());
m_cache.clear();
boost::property_tree::ptree pt;
try {
read_json(m_file_name, pt);
}
catch(boost::property_tree::json_parser::json_parser_error &e){
// no saved cache file, ignore
return;
}
std::string stored_device;
try {
stored_device = pt.get<std::string>("header.device");
}
catch(boost::property_tree::ptree_bad_path&){
return;
}
std::string stored_version;
try {
stored_version = pt.get<std::string>("header.version");
}
catch(boost::property_tree::ptree_bad_path&){
return;
}
if(stored_device == m_device_name && stored_version == version_string()){
typedef boost::property_tree::ptree::const_iterator pt_iter;
for(pt_iter iter = pt.begin(); iter != pt.end(); ++iter){
if(iter->first == "header"){
// skip header
continue;
}
boost::property_tree::ptree child_pt = pt.get_child(iter->first);
for(pt_iter child_iter = child_pt.begin(); child_iter != child_pt.end(); ++child_iter){
set(iter->first, child_iter->first, boost::lexical_cast<uint_>(child_iter->second.data()));
}
}
}
m_dirty = false;
}
#endif // BOOST_COMPUTE_USE_OFFLINE_CACHE
private:
bool m_dirty;
std::string m_device_name;
std::string m_file_name;
std::map<std::pair<std::string, std::string>, uint_> m_cache;
};
} // end detail namespace
} // end compute namespace
} // end boost namespace
#endif // BOOST_COMPUTE_DETAIL_PARAMETER_CACHE_HPP

73
include/boost/compute/detail/path.hpp Normal file
View File

@@ -0,0 +1,73 @@
//---------------------------------------------------------------------------//
// Copyright (c) 2013-2014 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_DETAIL_PATH_HPP
#define BOOST_COMPUTE_DETAIL_PATH_HPP
#include <boost/filesystem/path.hpp>
#include <boost/filesystem/operations.hpp>
#include <boost/compute/detail/getenv.hpp>
namespace boost {
namespace compute {
namespace detail {
// Path delimiter symbol for the current OS.
static const std::string& path_delim()
{
static const std::string delim =
boost::filesystem::path("/").make_preferred().string();
return delim;
}
// Path to appdata folder.
inline const std::string& appdata_path()
{
#ifdef WIN32
static const std::string appdata = detail::getenv("APPDATA")
+ path_delim() + "boost_compute";
#else
static const std::string appdata = detail::getenv("HOME")
+ path_delim() + ".boost_compute";
#endif
return appdata;
}
// Path to cached binaries.
inline std::string program_binary_path(const std::string &hash, bool create = false)
{
std::string dir = detail::appdata_path() + path_delim()
+ hash.substr(0, 2) + path_delim()
+ hash.substr(2);
if(create && !boost::filesystem::exists(dir)){
boost::filesystem::create_directories(dir);
}
return dir + path_delim();
}
// Path to parameter caches.
inline std::string parameter_cache_path(bool create = false)
{
const static std::string dir = appdata_path() + path_delim() + "tune";
if(create && !boost::filesystem::exists(dir)){
boost::filesystem::create_directories(dir);
}
return dir + path_delim();
}
} // end detail namespace
} // end compute namespace
} // end boost namespace
#endif // BOOST_COMPUTE_DETAIL_PATH_HPP

8
include/boost/compute/device.hpp
View File

@@ -11,11 +11,12 @@
#ifndef BOOST_COMPUTE_DEVICE_HPP
#define BOOST_COMPUTE_DEVICE_HPP
#include <algorithm>
#include <string>
#include <vector>
#include <boost/range/algorithm.hpp>
#include <boost/algorithm/string.hpp>
#include <boost/algorithm/string/split.hpp>
#include <boost/algorithm/string/classification.hpp>
#include <boost/compute/config.hpp>
#include <boost/compute/exception.hpp>
@@ -213,7 +214,8 @@ public:
{
const std::vector<std::string> extensions = this->extensions();
return boost::find(extensions, name) != extensions.end();
return std::find(
extensions.begin(), extensions.end(), name) != extensions.end();
}
/// Returns the number of address bits.

63
include/boost/compute/experimental/transform_if.hpp
View File

@@ -1,63 +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_EXPERIMENTAL_TRANSFORM_IF_HPP
#define BOOST_COMPUTE_EXPERIMENTAL_TRANSFORM_IF_HPP
#include <boost/compute/command_queue.hpp>
#include <boost/compute/detail/meta_kernel.hpp>
#include <boost/compute/detail/iterator_range_size.hpp>
#include <boost/compute/algorithm/detail/copy_on_device.hpp>
namespace boost {
namespace compute {
namespace experimental {
template<class InputIterator,
class OutputIterator,
class UnaryOperator,
class Predicate>
inline OutputIterator transform_if(InputIterator first,
InputIterator last,
OutputIterator result,
UnaryOperator op,
Predicate predicate,
command_queue &queue)
{
typedef typename
std::iterator_traits<InputIterator>::difference_type difference_type;
difference_type count = std::distance(first, last);
if(count < 1){
return result;
}
detail::meta_kernel k("transform_if");
k <<
k.if_(predicate(first[k.get_global_id(0)])) << "{\n" <<
result[k.get_global_id(0)] << '=' <<
op(first[k.get_global_id(0)]) << ";\n"
"}\n";
const device &device = queue.get_device();
const size_t work_group_size =
detail::pick_copy_work_group_size(count, device);
k.exec_1d(queue, 0, count, work_group_size);
return result + count;
}
} // end experimental namespace
} // end compute namespace
} // end boost namespace
#endif // BOOST_COMPUTE_ALGORITHM_TRANSFORM_IF_HPP

2
include/boost/compute/kernel.hpp
View File

@@ -197,7 +197,7 @@ public:
///
/// \see_opencl_ref{clGetKernelWorkGroupInfo}
template<class T>
T get_work_group_info(const device &device, cl_kernel_work_group_info info)
T get_work_group_info(const device &device, cl_kernel_work_group_info info) const
{
return detail::get_object_info<T>(clGetKernelWorkGroupInfo, m_kernel, info, device.id());
}

8
include/boost/compute/platform.hpp
View File

@@ -11,11 +11,12 @@
#ifndef BOOST_COMPUTE_PLATFORM_HPP
#define BOOST_COMPUTE_PLATFORM_HPP
#include <algorithm>
#include <string>
#include <vector>
#include <boost/range/algorithm.hpp>
#include <boost/algorithm/string.hpp>
#include <boost/algorithm/string/split.hpp>
#include <boost/algorithm/string/classification.hpp>
#include <boost/compute/cl.hpp>
#include <boost/compute/device.hpp>
@@ -112,7 +113,8 @@ public:
{
const std::vector<std::string> extensions = this->extensions();
return boost::find(extensions, name) != extensions.end();
return std::find(
extensions.begin(), extensions.end(), name) != extensions.end();
}
/// Returns a list of devices on the platform.

37
include/boost/compute/program.hpp
View File

@@ -28,9 +28,9 @@
#ifdef BOOST_COMPUTE_USE_OFFLINE_CACHE
#include <sstream>
#include <boost/optional.hpp>
#include <boost/filesystem.hpp>
#include <boost/compute/platform.hpp>
#include <boost/compute/detail/getenv.hpp>
#include <boost/compute/detail/path.hpp>
#include <boost/compute/detail/sha1.hpp>
#endif
@@ -558,41 +558,10 @@ public:
private:
#ifdef BOOST_COMPUTE_USE_OFFLINE_CACHE
// Path delimiter symbol for the current OS.
static const std::string& path_delim() {
static const std::string delim =
boost::filesystem::path("/").make_preferred().string();
return delim;
}
// Path to appdata folder.
static inline const std::string& appdata_path() {
#ifdef WIN32
static const std::string appdata = detail::getenv("APPDATA")
+ path_delim() + "boost_compute";
#else
static const std::string appdata = detail::getenv("HOME")
+ path_delim() + ".boost_compute";
#endif
return appdata;
}
// Path to cached binaries.
static std::string program_binary_path(const std::string &hash, bool create = false)
{
std::string dir = appdata_path() + path_delim()
+ hash.substr(0, 2) + path_delim()
+ hash.substr(2);
if (create) boost::filesystem::create_directories(dir);
return dir + path_delim();
}
// Saves program binaries for future reuse.
static void save_program_binary(const std::string &hash, const program &prog)
{
std::string fname = program_binary_path(hash, true) + "kernel";
std::string fname = detail::program_binary_path(hash, true) + "kernel";
std::ofstream bfile(fname.c_str(), std::ios::binary);
if (!bfile) return;
@@ -608,7 +577,7 @@ private:
const std::string &hash, const context &ctx
)
{
std::string fname = program_binary_path(hash) + "kernel";
std::string fname = detail::program_binary_path(hash) + "kernel";
std::ifstream bfile(fname.c_str(), std::ios::binary);
if (!bfile) return boost::optional<program>();

33
include/boost/compute/system.hpp
View File

@@ -15,7 +15,6 @@
#include <vector>
#include <cstdlib>
#include <boost/foreach.hpp>
#include <boost/throw_exception.hpp>
#include <boost/compute/cl.hpp>
@@ -85,7 +84,10 @@ public:
/// \throws no_device_found if no device with \p name is found.
static device find_device(const std::string &name)
{
BOOST_FOREACH(const device &device, devices()){
const std::vector<device> devices = system::devices();
for(size_t i = 0; i < devices.size(); i++){
const device& device = devices[i];
if(device.name() == name){
return device;
}
@@ -108,10 +110,13 @@ public:
{
std::vector<device> devices;
BOOST_FOREACH(const platform &platform, platforms()){
BOOST_FOREACH(const device &device, platform.devices()){
devices.push_back(device);
}
const std::vector<platform> platforms = system::platforms();
for(size_t i = 0; i < platforms.size(); i++){
const std::vector<device> platform_devices = platforms[i].devices();
devices.insert(
devices.end(), platform_devices.begin(), platform_devices.end()
);
}
return devices;
@@ -122,8 +127,9 @@ public:
{
size_t count = 0;
BOOST_FOREACH(const platform &platform, platforms()){
count += platform.device_count();
const std::vector<platform> platforms = system::platforms();
for(size_t i = 0; i < platforms.size(); i++){
count += platforms[i].device_count();
}
return count;
@@ -214,7 +220,8 @@ private:
const char *vendor = detail::getenv("BOOST_COMPUTE_DEFAULT_VENDOR");
if(name || type || platform || vendor){
BOOST_FOREACH(const device &device, devices_){
for(size_t i = 0; i < devices_.size(); i++){
const device& device = devices_[i];
if (name && !matches(device.name(), name))
continue;
@@ -237,14 +244,18 @@ private:
}
// find the first gpu device
BOOST_FOREACH(const device &device, devices_){
for(size_t i = 0; i < devices_.size(); i++){
const device& device = devices_[i];
if(device.type() == device::gpu){
return device;
}
}
// find the first cpu device
BOOST_FOREACH(const device &device, devices_){
for(size_t i = 0; i < devices_.size(); i++){
const device& device = devices_[i];
if(device.type() == device::cpu){
return device;
}

6
include/boost/compute/type_traits/is_fundamental.hpp
View File

@@ -33,7 +33,11 @@ struct is_fundamental : public boost::false_type {};
template<> struct is_fundamental<BOOST_PP_CAT(BOOST_PP_CAT(type, 2), _)> : boost::true_type {}; \
template<> struct is_fundamental<BOOST_PP_CAT(BOOST_PP_CAT(type, 4), _)> : boost::true_type {}; \
template<> struct is_fundamental<BOOST_PP_CAT(BOOST_PP_CAT(type, 8), _)> : boost::true_type {}; \
template<> struct is_fundamental<BOOST_PP_CAT(BOOST_PP_CAT(type, 16), _)> : boost::true_type {};
template<> struct is_fundamental<BOOST_PP_CAT(BOOST_PP_CAT(type, 16), _)> : boost::true_type {}; \
template<> struct is_fundamental<BOOST_PP_CAT(cl_, BOOST_PP_CAT(type, 2))> : boost::true_type {}; \
template<> struct is_fundamental<BOOST_PP_CAT(cl_, BOOST_PP_CAT(type, 4))> : boost::true_type {}; \
template<> struct is_fundamental<BOOST_PP_CAT(cl_, BOOST_PP_CAT(type, 8))> : boost::true_type {}; \
template<> struct is_fundamental<BOOST_PP_CAT(cl_, BOOST_PP_CAT(type, 16))> : boost::true_type {};
BOOST_COMPUTE_DETAIL_DECLARE_FUNDAMENTAL(char)
BOOST_COMPUTE_DETAIL_DECLARE_FUNDAMENTAL(uchar)

5
include/boost/compute/types/fundamental.hpp
View File

@@ -14,7 +14,10 @@
#include <cstring>
#include <ostream>
#include <boost/preprocessor.hpp>
#include <boost/preprocessor/cat.hpp>
#include <boost/preprocessor/comma.hpp>
#include <boost/preprocessor/repetition.hpp>
#include <boost/preprocessor/stringize.hpp>
#include <boost/compute/cl.hpp>

3
include/boost/compute/types/struct.hpp
View File

@@ -15,9 +15,10 @@
#include <boost/static_assert.hpp>
#include <boost/preprocessor/expr_if.hpp>
#include <boost/preprocessor/stringize.hpp>
#include <boost/preprocessor/seq/for_each.hpp>
#include <boost/preprocessor/seq/fold_left.hpp>
#include <boost/preprocessor/seq/for_each.hpp>
#include <boost/preprocessor/seq/transform.hpp>
#include <boost/compute/type_traits/type_definition.hpp>

5
include/boost/compute/types/tuple.hpp
View File

@@ -14,6 +14,9 @@
#include <string>
#include <utility>
#include <boost/preprocessor/enum.hpp>
#include <boost/preprocessor/expr_if.hpp>
#include <boost/preprocessor/repetition.hpp>
#include <boost/tuple/tuple.hpp>
#include <boost/compute/config.hpp>
@@ -25,8 +28,6 @@
#include <tuple>
#endif
#include <boost/preprocessor/repetition.hpp>
namespace boost {
namespace compute {
namespace detail {