diff options
Diffstat (limited to 'boost/compute/detail')
26 files changed, 3121 insertions, 0 deletions
diff --git a/boost/compute/detail/assert_cl_success.hpp b/boost/compute/detail/assert_cl_success.hpp new file mode 100644 index 0000000000..78acaf6caf --- /dev/null +++ b/boost/compute/detail/assert_cl_success.hpp @@ -0,0 +1,24 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_ASSERT_CL_SUCCESS_HPP +#define BOOST_COMPUTE_DETAIL_ASSERT_CL_SUCCESS_HPP + +#include <boost/assert.hpp> + +#if defined(BOOST_DISABLE_ASSERTS) || defined(NDEBUG) +#define BOOST_COMPUTE_ASSERT_CL_SUCCESS(function) \ + function +#else +#define BOOST_COMPUTE_ASSERT_CL_SUCCESS(function) \ + BOOST_ASSERT(function == CL_SUCCESS) +#endif + +#endif // BOOST_COMPUTE_DETAIL_ASSERT_CL_SUCCESS_HPP diff --git a/boost/compute/detail/buffer_value.hpp b/boost/compute/detail/buffer_value.hpp new file mode 100644 index 0000000000..6a4e78fc19 --- /dev/null +++ b/boost/compute/detail/buffer_value.hpp @@ -0,0 +1,178 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_BUFFER_VALUE_HPP +#define BOOST_COMPUTE_DETAIL_BUFFER_VALUE_HPP + +#include <boost/compute/context.hpp> +#include <boost/compute/command_queue.hpp> +#include <boost/compute/detail/device_ptr.hpp> +#include <boost/compute/detail/read_write_single_value.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class T> +class buffer_value +{ +public: + typedef T value_type; + + buffer_value() + { + } + + buffer_value(const value_type &value) + : m_value(value) + { + } + + // creates a reference for the value in buffer at index (in bytes). + buffer_value(const buffer &buffer, size_t index) + : m_buffer(buffer.get(), false), + m_index(index) + { + } + + buffer_value(const buffer_value<T> &other) + : m_buffer(other.m_buffer.get(), false), + m_index(other.m_index) + { + } + + ~buffer_value() + { + // set buffer to null so that its reference count will + // not be decremented when its destructor is called + m_buffer.get() = 0; + } + + operator value_type() const + { + if(m_buffer.get()){ + const context &context = m_buffer.get_context(); + const device &device = context.get_device(); + command_queue queue(context, device); + + return detail::read_single_value<T>(m_buffer, m_index / sizeof(T), queue); + } + else { + return m_value; + } + } + + buffer_value<T> operator-() const + { + return -T(*this); + } + + bool operator<(const T &value) const + { + return T(*this) < value; + } + + bool operator>(const T &value) const + { + return T(*this) > value; + } + + bool operator<=(const T &value) const + { + return T(*this) <= value; + } + + bool operator>=(const T &value) const + { + return T(*this) <= value; + } + + bool operator==(const T &value) const + { + return T(*this) == value; + } + + bool operator==(const buffer_value<T> &other) const + { + if(m_buffer.get() != other.m_buffer.get()){ + return false; + } + + if(m_buffer.get()){ + return m_index == other.m_index; + } + else { + return m_value == other.m_value; + } + } + + bool operator!=(const T &value) const + { + return T(*this) != value; + } + + buffer_value<T>& operator=(const T &value) + { + if(m_buffer.get()){ + const context &context = m_buffer.get_context(); + command_queue queue(context, context.get_device()); + + detail::write_single_value<T>(value, m_buffer, m_index / sizeof(T), queue); + + return *this; + } + else { + m_value = value; + return *this; + } + } + + buffer_value<T>& operator=(const buffer_value<T> &value) + { + return operator=(T(value)); + } + + detail::device_ptr<T> operator&() const + { + return detail::device_ptr<T>(m_buffer, m_index); + } + + buffer_value<T>& operator++() + { + if(m_buffer.get()){ + T value = T(*this); + value++; + *this = value; + } + else { + m_value++; + } + + return *this; + } + + buffer_value<T> operator++(int) + { + buffer_value<T> result(*this); + ++(*this); + return result; + } + +private: + const buffer m_buffer; + size_t m_index; + value_type m_value; +}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_BUFFER_VALUE_HPP diff --git a/boost/compute/detail/device_ptr.hpp b/boost/compute/detail/device_ptr.hpp new file mode 100644 index 0000000000..29ecd13631 --- /dev/null +++ b/boost/compute/detail/device_ptr.hpp @@ -0,0 +1,215 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DEVICE_PTR_HPP +#define BOOST_COMPUTE_DEVICE_PTR_HPP + +#include <boost/type_traits.hpp> +#include <boost/static_assert.hpp> + +#include <boost/compute/buffer.hpp> +#include <boost/compute/config.hpp> +#include <boost/compute/detail/is_buffer_iterator.hpp> +#include <boost/compute/detail/read_write_single_value.hpp> +#include <boost/compute/type_traits/is_device_iterator.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class T, class IndexExpr> +struct device_ptr_index_expr +{ + typedef T result_type; + + device_ptr_index_expr(const buffer &buffer, + uint_ index, + const IndexExpr &expr) + : m_buffer(buffer), + m_index(index), + m_expr(expr) + { + } + + operator T() const + { + BOOST_STATIC_ASSERT_MSG(boost::is_integral<IndexExpr>::value, + "Index expression must be integral"); + + BOOST_ASSERT(m_buffer.get()); + + const context &context = m_buffer.get_context(); + const device &device = context.get_device(); + command_queue queue(context, device); + + return detail::read_single_value<T>(m_buffer, m_expr, queue); + } + + const buffer &m_buffer; + uint_ m_index; + IndexExpr m_expr; +}; + +template<class T> +class device_ptr +{ +public: + typedef T value_type; + typedef std::size_t size_type; + typedef std::ptrdiff_t difference_type; + typedef std::random_access_iterator_tag iterator_category; + typedef T* pointer; + typedef T& reference; + + device_ptr() + : m_index(0) + { + } + + device_ptr(const buffer &buffer, size_t index = 0) + : m_buffer(buffer.get(), false), + m_index(index) + { + } + + device_ptr(const device_ptr<T> &other) + : m_buffer(other.m_buffer.get(), false), + m_index(other.m_index) + { + } + + device_ptr<T>& operator=(const device_ptr<T> &other) + { + if(this != &other){ + m_buffer.get() = other.m_buffer.get(); + m_index = other.m_index; + } + + return *this; + } + + #ifndef BOOST_COMPUTE_NO_RVALUE_REFERENCES + device_ptr(device_ptr<T>&& other) BOOST_NOEXCEPT + : m_buffer(other.m_buffer.get(), false), + m_index(other.m_index) + { + other.m_buffer.get() = 0; + } + + device_ptr<T>& operator=(device_ptr<T>&& other) BOOST_NOEXCEPT + { + m_buffer.get() = other.m_buffer.get(); + m_index = other.m_index; + + other.m_buffer.get() = 0; + + return *this; + } + #endif // BOOST_COMPUTE_NO_RVALUE_REFERENCES + + ~device_ptr() + { + // set buffer to null so that its reference count will + // not be decremented when its destructor is called + m_buffer.get() = 0; + } + + size_type get_index() const + { + return m_index; + } + + const buffer& get_buffer() const + { + return m_buffer; + } + + template<class OT> + device_ptr<OT> cast() const + { + return device_ptr<OT>(m_buffer, m_index); + } + + device_ptr<T> operator+(difference_type n) const + { + return device_ptr<T>(m_buffer, m_index + n); + } + + device_ptr<T> operator+(const device_ptr<T> &other) const + { + return device_ptr<T>(m_buffer, m_index + other.m_index); + } + + device_ptr<T>& operator+=(difference_type n) + { + m_index += static_cast<size_t>(n); + return *this; + } + + difference_type operator-(const device_ptr<T> &other) const + { + return static_cast<difference_type>(m_index - other.m_index); + } + + device_ptr<T>& operator-=(difference_type n) + { + m_index -= n; + return *this; + } + + bool operator==(const device_ptr<T> &other) const + { + return m_buffer.get() == other.m_buffer.get() && + m_index == other.m_index; + } + + bool operator!=(const device_ptr<T> &other) const + { + return !(*this == other); + } + + template<class Expr> + detail::device_ptr_index_expr<T, Expr> + operator[](const Expr &expr) const + { + BOOST_ASSERT(m_buffer.get()); + + return detail::device_ptr_index_expr<T, Expr>(m_buffer, + uint_(m_index), + expr); + } + +private: + const buffer m_buffer; + size_t m_index; +}; + +// is_buffer_iterator specialization for device_ptr +template<class Iterator> +struct is_buffer_iterator< + Iterator, + typename boost::enable_if< + boost::is_same< + device_ptr<typename Iterator::value_type>, + typename boost::remove_const<Iterator>::type + > + >::type +> : public boost::true_type {}; + +} // end detail namespace + +// is_device_iterator specialization for device_ptr +template<class T> +struct is_device_iterator<detail::device_ptr<T> > : boost::true_type {}; + +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DEVICE_PTR_HPP diff --git a/boost/compute/detail/diagnostic.hpp b/boost/compute/detail/diagnostic.hpp new file mode 100644 index 0000000000..76a69f6570 --- /dev/null +++ b/boost/compute/detail/diagnostic.hpp @@ -0,0 +1,112 @@ +//---------------------------------------------------------------------------// +// Copyright (c) 2016 Jakub Szuppe <j.szuppe@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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_DIAGNOSTIC_HPP +#define BOOST_COMPUTE_DETAIL_DIAGNOSTIC_HPP + +// Macros for suppressing warnings for GCC version 4.6 or later. Usage: +// +// BOOST_COMPUTE_BOOST_COMPUTE_GCC_DIAG_OFF(sign-compare); +// if(a < b){ +// BOOST_COMPUTE_BOOST_COMPUTE_GCC_DIAG_ON(sign-compare); +// +// Source: https://svn.boost.org/trac/boost/wiki/Guidelines/WarningsGuidelines +#if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 402 +#define BOOST_COMPUTE_GCC_DIAG_STR(s) #s +#define BOOST_COMPUTE_GCC_DIAG_JOINSTR(x,y) BOOST_COMPUTE_GCC_DIAG_STR(x ## y) +# define BOOST_COMPUTE_GCC_DIAG_DO_PRAGMA(x) _Pragma (#x) +# define BOOST_COMPUTE_GCC_DIAG_PRAGMA(x) BOOST_COMPUTE_GCC_DIAG_DO_PRAGMA(GCC diagnostic x) +# if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 406 +# define BOOST_COMPUTE_GCC_DIAG_OFF(x) BOOST_COMPUTE_GCC_DIAG_PRAGMA(push) \ + BOOST_COMPUTE_GCC_DIAG_PRAGMA(ignored BOOST_COMPUTE_GCC_DIAG_JOINSTR(-W,x)) +# define BOOST_COMPUTE_GCC_DIAG_ON(x) BOOST_COMPUTE_GCC_DIAG_PRAGMA(pop) +# else +# define BOOST_COMPUTE_GCC_DIAG_OFF(x) \ + BOOST_COMPUTE_GCC_DIAG_PRAGMA(ignored BOOST_COMPUTE_GCC_DIAG_JOINSTR(-W,x)) +# define BOOST_COMPUTE_GCC_DIAG_ON(x) \ + BOOST_COMPUTE_GCC_DIAG_PRAGMA(warning BOOST_COMPUTE_GCC_DIAG_JOINSTR(-W,x)) +# endif +#else // Ensure these macros do nothing for other compilers. +# define BOOST_COMPUTE_GCC_DIAG_OFF(x) +# define BOOST_COMPUTE_GCC_DIAG_ON(x) +#endif + +// Macros for suppressing warnings for Clang. +// +// BOOST_COMPUTE_BOOST_COMPUTE_CLANG_DIAG_OFF(sign-compare); +// if(a < b){ +// BOOST_COMPUTE_BOOST_COMPUTE_CLANG_DIAG_ON(sign-compare); +// +// Source: https://svn.boost.org/trac/boost/wiki/Guidelines/WarningsGuidelines +#ifdef __clang__ +# define BOOST_COMPUTE_CLANG_DIAG_STR(s) # s +// stringize s to "no-sign-compare" +# define BOOST_COMPUTE_CLANG_DIAG_JOINSTR(x,y) BOOST_COMPUTE_CLANG_DIAG_STR(x ## y) +// join -W with no-unused-variable to "-Wno-sign-compare" +# define BOOST_COMPUTE_CLANG_DIAG_DO_PRAGMA(x) _Pragma (#x) +// _Pragma is unary operator #pragma ("") +# define BOOST_COMPUTE_CLANG_DIAG_PRAGMA(x) \ + BOOST_COMPUTE_CLANG_DIAG_DO_PRAGMA(clang diagnostic x) +# define BOOST_COMPUTE_CLANG_DIAG_OFF(x) BOOST_COMPUTE_CLANG_DIAG_PRAGMA(push) \ + BOOST_COMPUTE_CLANG_DIAG_PRAGMA(ignored BOOST_COMPUTE_CLANG_DIAG_JOINSTR(-W,x)) +// For example: #pragma clang diagnostic ignored "-Wno-sign-compare" +# define BOOST_COMPUTE_CLANG_DIAG_ON(x) BOOST_COMPUTE_CLANG_DIAG_PRAGMA(pop) +// For example: #pragma clang diagnostic warning "-Wno-sign-compare" +#else // Ensure these macros do nothing for other compilers. +# define BOOST_COMPUTE_CLANG_DIAG_OFF(x) +# define BOOST_COMPUTE_CLANG_DIAG_ON(x) +# define BOOST_COMPUTE_CLANG_DIAG_PRAGMA(x) +#endif + +// Macros for suppressing warnings for MSVC. Usage: +// +// BOOST_COMPUTE_BOOST_COMPUTE_MSVC_DIAG_OFF(4018); //sign-compare +// if(a < b){ +// BOOST_COMPUTE_BOOST_COMPUTE_MSVC_DIAG_ON(4018); +// +#if defined(_MSC_VER) +# define BOOST_COMPUTE_MSVC_DIAG_DO_PRAGMA(x) __pragma(x) +# define BOOST_COMPUTE_MSVC_DIAG_PRAGMA(x) \ + BOOST_COMPUTE_MSVC_DIAG_DO_PRAGMA(warning(x)) +# define BOOST_COMPUTE_MSVC_DIAG_OFF(x) BOOST_COMPUTE_MSVC_DIAG_PRAGMA(push) \ + BOOST_COMPUTE_MSVC_DIAG_PRAGMA(disable: x) +# define BOOST_COMPUTE_MSVC_DIAG_ON(x) BOOST_COMPUTE_MSVC_DIAG_PRAGMA(pop) +#else // Ensure these macros do nothing for other compilers. +# define BOOST_COMPUTE_MSVC_DIAG_OFF(x) +# define BOOST_COMPUTE_MSVC_DIAG_ON(x) +#endif + +// Macros for suppressing warnings for GCC, Clang and MSVC. Usage: +// +// BOOST_COMPUTE_DIAG_OFF(sign-compare, sign-compare, 4018); +// if(a < b){ +// BOOST_COMPUTE_DIAG_ON(sign-compare, sign-compare, 4018); +// +#if defined(_MSC_VER) // MSVC +# define BOOST_COMPUTE_DIAG_OFF(gcc, clang, msvc) BOOST_COMPUTE_MSVC_DIAG_OFF(msvc) +# define BOOST_COMPUTE_DIAG_ON(gcc, clang, msvc) BOOST_COMPUTE_MSVC_DIAG_ON(msvc) +#elif defined(__clang__) // Clang +# define BOOST_COMPUTE_DIAG_OFF(gcc, clang, msvc) BOOST_COMPUTE_CLANG_DIAG_OFF(clang) +# define BOOST_COMPUTE_DIAG_ON(gcc, clang, msvc) BOOST_COMPUTE_CLANG_DIAG_ON(clang) +#elif defined(__GNUC__) // GCC/G++ +# define BOOST_COMPUTE_DIAG_OFF(gcc, clang, msvc) BOOST_COMPUTE_GCC_DIAG_OFF(gcc) +# define BOOST_COMPUTE_DIAG_ON(gcc, clang, msvc) BOOST_COMPUTE_GCC_DIAG_ON(gcc) +#else // Ensure these macros do nothing for other compilers. +# define BOOST_COMPUTE_DIAG_OFF(gcc, clang, msvc) +# define BOOST_COMPUTE_DIAG_ON(gcc, clang, msvc) +#endif + +#define BOOST_COMPUTE_DISABLE_DEPRECATED_DECLARATIONS() \ + BOOST_COMPUTE_DIAG_OFF(deprecated-declarations, deprecated-declarations, 4996) +#define BOOST_COMPUTE_ENABLE_DEPRECATED_DECLARATIONS() \ + BOOST_COMPUTE_DIAG_ON(deprecated-declarations, deprecated-declarations, 4996); + + +#endif /* BOOST_COMPUTE_DETAIL_DIAGNOSTIC_HPP */ diff --git a/boost/compute/detail/duration.hpp b/boost/compute/detail/duration.hpp new file mode 100644 index 0000000000..601f12d291 --- /dev/null +++ b/boost/compute/detail/duration.hpp @@ -0,0 +1,50 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_DURATION_HPP +#define BOOST_COMPUTE_DETAIL_DURATION_HPP + +#include <boost/config.hpp> + +#ifndef BOOST_COMPUTE_NO_HDR_CHRONO +#include <chrono> +#endif + +#include <boost/chrono/duration.hpp> + +namespace boost { +namespace compute { +namespace detail { + +#ifndef BOOST_COMPUTE_NO_HDR_CHRONO +template<class Rep, class Period> +inline std::chrono::duration<Rep, Period> +make_duration_from_nanoseconds(std::chrono::duration<Rep, Period>, size_t nanoseconds) +{ + return std::chrono::duration_cast<std::chrono::duration<Rep, Period> >( + std::chrono::nanoseconds(nanoseconds) + ); +} +#endif // BOOST_COMPUTE_NO_HDR_CHRONO + +template<class Rep, class Period> +inline boost::chrono::duration<Rep, Period> +make_duration_from_nanoseconds(boost::chrono::duration<Rep, Period>, size_t nanoseconds) +{ + return boost::chrono::duration_cast<boost::chrono::duration<Rep, Period> >( + boost::chrono::nanoseconds(nanoseconds) + ); +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_DURATION_HPP diff --git a/boost/compute/detail/get_object_info.hpp b/boost/compute/detail/get_object_info.hpp new file mode 100644 index 0000000000..cdc20cbc13 --- /dev/null +++ b/boost/compute/detail/get_object_info.hpp @@ -0,0 +1,216 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_GET_OBJECT_INFO_HPP +#define BOOST_COMPUTE_DETAIL_GET_OBJECT_INFO_HPP + +#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> +#include <boost/compute/exception/opencl_error.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class Function, class Object, class AuxInfo> +struct bound_info_function +{ + bound_info_function(Function function, Object object, AuxInfo aux_info) + : m_function(function), + m_object(object), + m_aux_info(aux_info) + { + } + + template<class Info> + cl_int operator()(Info info, size_t size, void *value, size_t *size_ret) const + { + return m_function(m_object, m_aux_info, info, size, value, size_ret); + } + + Function m_function; + Object m_object; + AuxInfo m_aux_info; +}; + +template<class Function, class Object> +struct bound_info_function<Function, Object, void> +{ + bound_info_function(Function function, Object object) + : m_function(function), + m_object(object) + { + } + + template<class Info> + cl_int operator()(Info info, size_t size, void *value, size_t *size_ret) const + { + return m_function(m_object, info, size, value, size_ret); + } + + Function m_function; + Object m_object; +}; + +template<class Function, class Object> +inline bound_info_function<Function, Object, void> +bind_info_function(Function f, Object o) +{ + return bound_info_function<Function, Object, void>(f, o); +} + +template<class Function, class Object, class AuxInfo> +inline bound_info_function<Function, Object, AuxInfo> +bind_info_function(Function f, Object o, AuxInfo j) +{ + return bound_info_function<Function, Object, AuxInfo>(f, o, j); +} + +// default implementation +template<class T> +struct get_object_info_impl +{ + template<class Function, class Info> + T operator()(Function function, Info info) const + { + T value; + + cl_int ret = function(info, sizeof(T), &value, 0); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(opencl_error(ret)); + } + + return value; + } +}; + +// specialization for bool +template<> +struct get_object_info_impl<bool> +{ + template<class Function, class Info> + bool operator()(Function function, Info info) const + { + cl_bool value; + + cl_int ret = function(info, sizeof(cl_bool), &value, 0); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(opencl_error(ret)); + } + + return value == CL_TRUE; + } +}; + +// specialization for std::string +template<> +struct get_object_info_impl<std::string> +{ + template<class Function, class Info> + std::string operator()(Function function, Info info) const + { + size_t size = 0; + + cl_int ret = function(info, 0, 0, &size); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(opencl_error(ret)); + } + + if(size == 0){ + return std::string(); + } + + std::string value(size - 1, 0); + + ret = function(info, size, &value[0], 0); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(opencl_error(ret)); + } + + return value; + } +}; + +// specialization for std::vector<T> +template<class T> +struct get_object_info_impl<std::vector<T> > +{ + template<class Function, class Info> + std::vector<T> operator()(Function function, Info info) const + { + size_t size = 0; + + cl_int ret = function(info, 0, 0, &size); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(opencl_error(ret)); + } + + std::vector<T> vector(size / sizeof(T)); + ret = function(info, size, &vector[0], 0); + if(ret != CL_SUCCESS){ + BOOST_THROW_EXCEPTION(opencl_error(ret)); + } + + return vector; + } +}; + +// returns the value (of type T) from the given clGet*Info() function call. +template<class T, class Function, class Object, class Info> +inline T get_object_info(Function f, Object o, Info i) +{ + return get_object_info_impl<T>()(bind_info_function(f, o), i); +} + +template<class T, class Function, class Object, class Info, class AuxInfo> +inline T get_object_info(Function f, Object o, Info i, AuxInfo j) +{ + return get_object_info_impl<T>()(bind_info_function(f, o, j), i); +} + +// returns the value type for the clGet*Info() call on Object with Enum. +template<class Object, int Enum> +struct get_object_info_type; + +// defines the object::get_info<Enum>() specialization +#define BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATION(object_type, result_type, value) \ + namespace detail { \ + template<> struct get_object_info_type<object_type, value> { typedef result_type type; }; \ + } \ + template<> inline result_type object_type::get_info<value>() const \ + { \ + return get_info<result_type>(value); \ + } + +// used by BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATIONS() +#define BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_IMPL(r, data, elem) \ + BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATION( \ + data, BOOST_PP_TUPLE_ELEM(2, 0, elem), BOOST_PP_TUPLE_ELEM(2, 1, elem) \ + ) + +// defines the object::get_info<Enum>() specialization for each +// (result_type, value) tuple in seq for object_type. +#define BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_SPECIALIZATIONS(object_type, seq) \ + BOOST_PP_SEQ_FOR_EACH( \ + BOOST_COMPUTE_DETAIL_DEFINE_GET_INFO_IMPL, object_type, seq \ + ) + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_GET_OBJECT_INFO_HPP diff --git a/boost/compute/detail/getenv.hpp b/boost/compute/detail/getenv.hpp new file mode 100644 index 0000000000..ceb3605d5a --- /dev/null +++ b/boost/compute/detail/getenv.hpp @@ -0,0 +1,36 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_GETENV_HPP +#define BOOST_COMPUTE_DETAIL_GETENV_HPP + +#include <cstdlib> + +namespace boost { +namespace compute { +namespace detail { + +inline const char* getenv(const char *env_var) +{ +#ifdef _MSC_VER +# pragma warning(push) +# pragma warning(disable: 4996) +#endif + return std::getenv(env_var); +#ifdef _MSC_VER +# pragma warning(pop) +#endif +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_GETENV_HPP diff --git a/boost/compute/detail/global_static.hpp b/boost/compute/detail/global_static.hpp new file mode 100644 index 0000000000..d8014e4252 --- /dev/null +++ b/boost/compute/detail/global_static.hpp @@ -0,0 +1,37 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_GLOBAL_STATIC_HPP +#define BOOST_COMPUTE_DETAIL_GLOBAL_STATIC_HPP + +#include <boost/compute/config.hpp> + +#ifdef BOOST_COMPUTE_THREAD_SAFE +# ifdef BOOST_COMPUTE_HAVE_THREAD_LOCAL + // use c++11 thread local storage +# define BOOST_COMPUTE_DETAIL_GLOBAL_STATIC(type, name, ctor) \ + thread_local type name ctor; +# else + // use thread_specific_ptr from boost.thread +# include <boost/thread/tss.hpp> +# define BOOST_COMPUTE_DETAIL_GLOBAL_STATIC(type, name, ctor) \ + static ::boost::thread_specific_ptr< type > BOOST_PP_CAT(name, _tls_ptr_); \ + if(!BOOST_PP_CAT(name, _tls_ptr_).get()){ \ + BOOST_PP_CAT(name, _tls_ptr_).reset(new type ctor); \ + } \ + static type &name = *BOOST_PP_CAT(name, _tls_ptr_); +# endif +#else + // no thread-safety, just use static +# define BOOST_COMPUTE_DETAIL_GLOBAL_STATIC(type, name, ctor) \ + static type name ctor; +#endif + +#endif // BOOST_COMPUTE_DETAIL_GLOBAL_STATIC_HPP diff --git a/boost/compute/detail/is_buffer_iterator.hpp b/boost/compute/detail/is_buffer_iterator.hpp new file mode 100644 index 0000000000..c0caa050d6 --- /dev/null +++ b/boost/compute/detail/is_buffer_iterator.hpp @@ -0,0 +1,30 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_IS_BUFFER_ITERATOR_HPP +#define BOOST_COMPUTE_DETAIL_IS_BUFFER_ITERATOR_HPP + +#include <boost/config.hpp> +#include <boost/type_traits.hpp> +#include <boost/utility/enable_if.hpp> + +namespace boost { +namespace compute { +namespace detail { + +// default = false +template<class Iterator, class Enable = void> +struct is_buffer_iterator : public boost::false_type {}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_IS_BUFFER_ITERATOR_HPP diff --git a/boost/compute/detail/is_contiguous_iterator.hpp b/boost/compute/detail/is_contiguous_iterator.hpp new file mode 100644 index 0000000000..d0889b2f9e --- /dev/null +++ b/boost/compute/detail/is_contiguous_iterator.hpp @@ -0,0 +1,118 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_IS_CONTIGUOUS_ITERATOR_HPP +#define BOOST_COMPUTE_DETAIL_IS_CONTIGUOUS_ITERATOR_HPP + +#include <vector> +#include <valarray> + +#include <boost/config.hpp> +#include <boost/type_traits.hpp> +#include <boost/utility/enable_if.hpp> + +namespace boost { +namespace compute { +namespace detail { + +// default = false +template<class Iterator, class Enable = void> +struct _is_contiguous_iterator : public boost::false_type {}; + +// std::vector<T>::iterator = true +template<class Iterator> +struct _is_contiguous_iterator< + Iterator, + typename boost::enable_if< + typename boost::is_same< + Iterator, + typename std::vector<typename Iterator::value_type>::iterator + >::type + >::type +> : public boost::true_type {}; + +// std::vector<T>::const_iterator = true +template<class Iterator> +struct _is_contiguous_iterator< + Iterator, + typename boost::enable_if< + typename boost::is_same< + Iterator, + typename std::vector<typename Iterator::value_type>::const_iterator + >::type + >::type +> : public boost::true_type {}; + +// std::valarray<T>::iterator = true +template<class Iterator> +struct _is_contiguous_iterator< + Iterator, + typename boost::enable_if< + typename boost::is_same< + Iterator, + typename std::valarray<typename Iterator::value_type>::iterator + >::type + >::type +> : public boost::true_type {}; + +// std::valarray<T>::const_iterator = true +template<class Iterator> +struct _is_contiguous_iterator< + Iterator, + typename boost::enable_if< + typename boost::is_same< + Iterator, + typename std::valarray<typename Iterator::value_type>::const_iterator + >::type + >::type +> : public boost::true_type {}; + +// T* = true +template<class Iterator> +struct _is_contiguous_iterator< + Iterator, + typename boost::enable_if< + boost::is_pointer<Iterator> + >::type +> : public boost::true_type {}; + +// the is_contiguous_iterator meta-function returns true if Iterator points +// to a range of contiguous values. examples of contiguous iterators are +// std::vector<>::iterator and float*. examples of non-contiguous iterators +// are std::set<>::iterator and std::insert_iterator<>. +// +// the implementation consists of two phases. the first checks that value_type +// for the iterator is not void. this must be done as for many containers void +// is not a valid value_type (ex. std::vector<void>::iterator is not valid). +// after ensuring a non-void value_type, the _is_contiguous_iterator function +// is invoked. it has specializations retuning true for all (known) contiguous +// iterators types and a default value of false. +template<class Iterator, class Enable = void> +struct is_contiguous_iterator : + public _is_contiguous_iterator< + typename boost::remove_cv<Iterator>::type + > {}; + +// value_type of void = false +template<class Iterator> +struct is_contiguous_iterator< + Iterator, + typename boost::enable_if< + typename boost::is_void< + typename Iterator::value_type + >::type + >::type +> : public boost::false_type {}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_IS_CONTIGUOUS_ITERATOR_HPP diff --git a/boost/compute/detail/iterator_plus_distance.hpp b/boost/compute/detail/iterator_plus_distance.hpp new file mode 100644 index 0000000000..26e95f16c0 --- /dev/null +++ b/boost/compute/detail/iterator_plus_distance.hpp @@ -0,0 +1,53 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_ITERATOR_PLUS_DISTANCE_HPP +#define BOOST_COMPUTE_DETAIL_ITERATOR_PLUS_DISTANCE_HPP + +#include <iterator> + +namespace boost { +namespace compute { +namespace detail { + +template<class Iterator, class Distance, class Tag> +inline Iterator iterator_plus_distance(Iterator i, Distance n, Tag) +{ + while(n--){ i++; } + + return i; +} + +template<class Iterator, class Distance> +inline Iterator iterator_plus_distance(Iterator i, + Distance n, + std::random_access_iterator_tag) +{ + typedef typename + std::iterator_traits<Iterator>::difference_type difference_type; + + return i + static_cast<difference_type>(n); +} + +// similar to std::advance() except returns the advanced iterator and +// also works with iterators that don't define difference_type +template<class Iterator, class Distance> +inline Iterator iterator_plus_distance(Iterator i, Distance n) +{ + typedef typename std::iterator_traits<Iterator>::iterator_category tag; + + return iterator_plus_distance(i, n, tag()); +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_ITERATOR_PLUS_DISTANCE_HPP diff --git a/boost/compute/detail/iterator_range_size.hpp b/boost/compute/detail/iterator_range_size.hpp new file mode 100644 index 0000000000..67a675f833 --- /dev/null +++ b/boost/compute/detail/iterator_range_size.hpp @@ -0,0 +1,44 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_ITERATOR_RANGE_SIZE_H +#define BOOST_COMPUTE_DETAIL_ITERATOR_RANGE_SIZE_H + +#include <cstddef> +#include <algorithm> +#include <iterator> + +namespace boost { +namespace compute { +namespace detail { + +// This is a convenience function which returns the size of a range +// bounded by two iterators. This function has two differences from +// the std::distance() function: 1) the return type (size_t) is +// unsigned, and 2) the return value is always positive. +template<class Iterator> +inline size_t iterator_range_size(Iterator first, Iterator last) +{ + typedef typename + std::iterator_traits<Iterator>::difference_type + difference_type; + + difference_type difference = std::distance(first, last); + + return static_cast<size_t>( + (std::max)(difference, static_cast<difference_type>(0)) + ); +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_ITERATOR_RANGE_SIZE_H diff --git a/boost/compute/detail/iterator_traits.hpp b/boost/compute/detail/iterator_traits.hpp new file mode 100644 index 0000000000..45f0f683e6 --- /dev/null +++ b/boost/compute/detail/iterator_traits.hpp @@ -0,0 +1,35 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_ITERATOR_TRAITS_HPP +#define BOOST_COMPUTE_DETAIL_ITERATOR_TRAITS_HPP + +#include <iterator> + +#include <boost/compute/detail/is_contiguous_iterator.hpp> +#include <boost/compute/type_traits/is_device_iterator.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class Iterator> +struct iterator_traits : public std::iterator_traits<Iterator> +{ + static const bool is_contiguous = is_contiguous_iterator<Iterator>::value; + static const bool is_on_device = is_device_iterator<Iterator>::value; + static const bool is_on_host = !is_on_device; +}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_ITERATOR_TRAITS_HPP diff --git a/boost/compute/detail/literal.hpp b/boost/compute/detail/literal.hpp new file mode 100644 index 0000000000..0d23b1d4d2 --- /dev/null +++ b/boost/compute/detail/literal.hpp @@ -0,0 +1,45 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_LITERAL_HPP +#define BOOST_COMPUTE_DETAIL_LITERAL_HPP + +#include <iomanip> +#include <limits> +#include <sstream> + +#include <boost/type_traits/is_same.hpp> + +#include <boost/compute/types/fundamental.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class T> +std::string make_literal(T x) +{ + std::stringstream s; + s << std::setprecision(std::numeric_limits<T>::digits10) + << std::scientific + << x; + + if(boost::is_same<T, float>::value || boost::is_same<T, float_>::value){ + s << "f"; + } + + return s.str(); +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_LITERAL_HPP diff --git a/boost/compute/detail/lru_cache.hpp b/boost/compute/detail/lru_cache.hpp new file mode 100644 index 0000000000..fe1a56f74b --- /dev/null +++ b/boost/compute/detail/lru_cache.hpp @@ -0,0 +1,139 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_LRU_CACHE_HPP +#define BOOST_COMPUTE_DETAIL_LRU_CACHE_HPP + +#include <map> +#include <list> +#include <utility> + +#include <boost/optional.hpp> + +namespace boost { +namespace compute { +namespace detail { + +// a cache which evicts the least recently used item when it is full +template<class Key, class Value> +class lru_cache +{ +public: + typedef Key key_type; + typedef Value value_type; + typedef std::list<key_type> list_type; + typedef std::map< + key_type, + std::pair<value_type, typename list_type::iterator> + > map_type; + + lru_cache(size_t capacity) + : m_capacity(capacity) + { + } + + ~lru_cache() + { + } + + size_t size() const + { + return m_map.size(); + } + + size_t capacity() const + { + return m_capacity; + } + + bool empty() const + { + return m_map.empty(); + } + + bool contains(const key_type &key) + { + return m_map.find(key) != m_map.end(); + } + + void insert(const key_type &key, const value_type &value) + { + typename map_type::iterator i = m_map.find(key); + if(i == m_map.end()){ + // insert item into the cache, but first check if it is full + if(size() >= m_capacity){ + // cache is full, evict the least recently used item + evict(); + } + + // insert the new item + m_list.push_front(key); + m_map[key] = std::make_pair(value, m_list.begin()); + } + } + + boost::optional<value_type> get(const key_type &key) + { + // lookup value in the cache + typename map_type::iterator i = m_map.find(key); + if(i == m_map.end()){ + // value not in cache + return boost::none; + } + + // return the value, but first update its place in the most + // recently used list + typename list_type::iterator j = i->second.second; + if(j != m_list.begin()){ + // move item to the front of the most recently used list + m_list.erase(j); + m_list.push_front(key); + + // update iterator in map + j = m_list.begin(); + const value_type &value = i->second.first; + m_map[key] = std::make_pair(value, j); + + // return the value + return value; + } + else { + // the item is already at the front of the most recently + // used list so just return it + return i->second.first; + } + } + + void clear() + { + m_map.clear(); + m_list.clear(); + } + +private: + void evict() + { + // evict item from the end of most recently used list + typename list_type::iterator i = --m_list.end(); + m_map.erase(*i); + m_list.erase(i); + } + +private: + map_type m_map; + list_type m_list; + size_t m_capacity; +}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_LRU_CACHE_HPP diff --git a/boost/compute/detail/meta_kernel.hpp b/boost/compute/detail/meta_kernel.hpp new file mode 100644 index 0000000000..7be778b025 --- /dev/null +++ b/boost/compute/detail/meta_kernel.hpp @@ -0,0 +1,1054 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_META_KERNEL_HPP +#define BOOST_COMPUTE_DETAIL_META_KERNEL_HPP + +#include <set> +#include <string> +#include <vector> +#include <iomanip> +#include <sstream> +#include <utility> + +#include <boost/tuple/tuple.hpp> +#include <boost/type_traits.hpp> +#include <boost/lexical_cast.hpp> +#include <boost/static_assert.hpp> +#include <boost/algorithm/string/find.hpp> +#include <boost/preprocessor/repetition.hpp> + +#include <boost/compute/kernel.hpp> +#include <boost/compute/closure.hpp> +#include <boost/compute/function.hpp> +#include <boost/compute/functional.hpp> +#include <boost/compute/type_traits.hpp> +#include <boost/compute/command_queue.hpp> +#include <boost/compute/image/image2d.hpp> +#include <boost/compute/image/image_sampler.hpp> +#include <boost/compute/memory_object.hpp> +#include <boost/compute/detail/device_ptr.hpp> +#include <boost/compute/detail/sha1.hpp> +#include <boost/compute/utility/program_cache.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class T> +class meta_kernel_variable +{ +public: + typedef T result_type; + + meta_kernel_variable(const std::string &name) + : m_name(name) + { + } + + meta_kernel_variable(const meta_kernel_variable &other) + : m_name(other.m_name) + { + } + + meta_kernel_variable& operator=(const meta_kernel_variable &other) + { + if(this != &other){ + m_name = other.m_name; + } + + return *this; + } + + ~meta_kernel_variable() + { + } + + std::string name() const + { + return m_name; + } + +private: + std::string m_name; +}; + +template<class T> +class meta_kernel_literal +{ +public: + typedef T result_type; + + meta_kernel_literal(const T &value) + : m_value(value) + { + } + + meta_kernel_literal(const meta_kernel_literal &other) + : m_value(other.m_value) + { + } + + meta_kernel_literal& operator=(const meta_kernel_literal &other) + { + if(this != &other){ + m_value = other.m_value; + } + + return *this; + } + + ~meta_kernel_literal() + { + } + + const T& value() const + { + return m_value; + } + +private: + T m_value; +}; + +struct meta_kernel_stored_arg +{ + meta_kernel_stored_arg() + : m_size(0), + m_value(0) + { + } + + meta_kernel_stored_arg(const meta_kernel_stored_arg &other) + : m_size(0), + m_value(0) + { + set_value(other.m_size, other.m_value); + } + + meta_kernel_stored_arg& operator=(const meta_kernel_stored_arg &other) + { + if(this != &other){ + set_value(other.m_size, other.m_value); + } + + return *this; + } + + template<class T> + meta_kernel_stored_arg(const T &value) + : m_size(0), + m_value(0) + { + set_value(value); + } + + ~meta_kernel_stored_arg() + { + if(m_value){ + std::free(m_value); + } + } + + void set_value(size_t size, const void *value) + { + if(m_value){ + std::free(m_value); + } + + m_size = size; + + if(value){ + m_value = std::malloc(size); + std::memcpy(m_value, value, size); + } + else { + m_value = 0; + } + } + + template<class T> + void set_value(const T &value) + { + set_value(sizeof(T), boost::addressof(value)); + } + + size_t m_size; + void *m_value; +}; + +struct meta_kernel_buffer_info +{ + meta_kernel_buffer_info(const buffer &buffer, + const std::string &id, + memory_object::address_space addr_space, + size_t i) + : m_mem(buffer.get()), + identifier(id), + address_space(addr_space), + index(i) + { + } + + cl_mem m_mem; + std::string identifier; + memory_object::address_space address_space; + size_t index; +}; + +class meta_kernel; + +template<class Type> +struct inject_type_impl +{ + void operator()(meta_kernel &) + { + // default implementation does nothing + } +}; + +#define BOOST_COMPUTE_META_KERNEL_DECLARE_SCALAR_TYPE_STREAM_OPERATOR(type) \ + meta_kernel& operator<<(const type &x) \ + { \ + m_source << x; \ + return *this; \ + } + +#define BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(type) \ + meta_kernel& operator<<(const type &x) \ + { \ + m_source << "(" << type_name<type>() << ")"; \ + m_source << "("; \ + for(size_t i = 0; i < vector_size<type>::value; i++){ \ + *this << lit(x[i]); \ + \ + if(i != vector_size<type>::value - 1){ \ + m_source << ","; \ + } \ + } \ + m_source << ")"; \ + return *this; \ + } + +#define BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(type) \ + BOOST_COMPUTE_META_KERNEL_DECLARE_SCALAR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(type, _)) \ + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 2), _)) \ + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 4), _)) \ + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 8), _)) \ + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 16), _)) + +class meta_kernel +{ +public: + template<class T> + class argument + { + public: + argument(const std::string &name, size_t index) + : m_name(name), + m_index(index) + { + } + + const std::string &name() const + { + return m_name; + } + + size_t index() const + { + return m_index; + } + + private: + std::string m_name; + size_t m_index; + }; + + explicit meta_kernel(const std::string &name) + : m_name(name) + { + } + + meta_kernel(const meta_kernel &other) + { + m_source.str(other.m_source.str()); + } + + meta_kernel& operator=(const meta_kernel &other) + { + if(this != &other){ + m_source.str(other.m_source.str()); + } + + return *this; + } + + ~meta_kernel() + { + } + + std::string name() const + { + return m_name; + } + + std::string source() const + { + std::stringstream stream; + + // add pragmas + if(!m_pragmas.empty()){ + stream << m_pragmas << "\n"; + } + + // add macros + stream << "#define boost_pair_type(t1, t2) _pair_ ## t1 ## _ ## t2 ## _t\n"; + stream << "#define boost_pair_get(x, n) (n == 0 ? x.first ## x.second)\n"; + stream << "#define boost_make_pair(t1, x, t2, y) (boost_pair_type(t1, t2)) { x, y }\n"; + stream << "#define boost_tuple_get(x, n) (x.v ## n)\n"; + + // add type declaration source + stream << m_type_declaration_source.str() << "\n"; + + // add external function source + stream << m_external_function_source.str() << "\n"; + + // add kernel source + stream << "__kernel void " << m_name + << "(" << boost::join(m_args, ", ") << ")\n" + << "{\n" << m_source.str() << "\n}\n"; + + return stream.str(); + } + + kernel compile(const context &context, const std::string &options = std::string()) + { + // generate the program source + std::string source = this->source(); + + // generate cache key + std::string cache_key = "__boost_meta_kernel_" + + static_cast<std::string>(detail::sha1(source)); + + // load program cache + boost::shared_ptr<program_cache> cache = + program_cache::get_global_cache(context); + + // load (or build) program from cache + ::boost::compute::program program = + cache->get_or_build(cache_key, options, source, context); + + // create kernel + ::boost::compute::kernel kernel = program.create_kernel(name()); + + // bind stored args + for(size_t i = 0; i < m_stored_args.size(); i++){ + const detail::meta_kernel_stored_arg &arg = m_stored_args[i]; + + if(arg.m_size != 0){ + kernel.set_arg(i, arg.m_size, arg.m_value); + } + } + + // bind buffer args + for(size_t i = 0; i < m_stored_buffers.size(); i++){ + const detail::meta_kernel_buffer_info &bi = m_stored_buffers[i]; + + kernel.set_arg(bi.index, bi.m_mem); + } + + return kernel; + } + + template<class T> + size_t add_arg(const std::string &name) + { + std::stringstream stream; + stream << type<T>() << " " << name; + + // add argument to list + m_args.push_back(stream.str()); + + // return index + return m_args.size() - 1; + } + + template<class T> + size_t add_arg(memory_object::address_space address_space, + const std::string &name) + { + return add_arg_with_qualifiers<T>(address_space_prefix(address_space), name); + } + + template<class T> + void set_arg(size_t index, const T &value) + { + if(index >= m_stored_args.size()){ + m_stored_args.resize(index + 1); + } + + m_stored_args[index] = detail::meta_kernel_stored_arg(value); + } + + void set_arg(size_t index, const memory_object &mem) + { + set_arg<cl_mem>(index, mem.get()); + } + + void set_arg(size_t index, const image_sampler &sampler) + { + set_arg<cl_sampler>(index, cl_sampler(sampler)); + } + + template<class T> + size_t add_set_arg(const std::string &name, const T &value) + { + size_t index = add_arg<T>(name); + set_arg<T>(index, value); + return index; + } + + void add_extension_pragma(const std::string &extension, + const std::string &value = "enable") + { + m_pragmas += "#pragma OPENCL EXTENSION " + extension + " : " + value + "\n"; + } + + void add_extension_pragma(const std::string &extension, + const std::string &value) const + { + return const_cast<meta_kernel *>(this)->add_extension_pragma(extension, value); + } + + template<class T> + std::string type() const + { + std::stringstream stream; + + // const qualifier + if(boost::is_const<T>::value){ + stream << "const "; + } + + // volatile qualifier + if(boost::is_volatile<T>::value){ + stream << "volatile "; + } + + // type + typedef + typename boost::remove_cv< + typename boost::remove_pointer<T>::type + >::type Type; + stream << type_name<Type>(); + + // pointer + if(boost::is_pointer<T>::value){ + stream << "*"; + } + + // inject type pragmas and/or definitions + inject_type<Type>(); + + return stream.str(); + } + + template<class T> + std::string decl(const std::string &name) const + { + return type<T>() + " " + name; + } + + template<class T, class Expr> + std::string decl(const std::string &name, const Expr &init) const + { + meta_kernel tmp((std::string())); + tmp << tmp.decl<T>(name) << " = " << init; + return tmp.m_source.str(); + } + + template<class T> + detail::meta_kernel_variable<T> var(const std::string &name) const + { + type<T>(); + + return make_var<T>(name); + } + + template<class T> + detail::meta_kernel_literal<T> lit(const T &value) const + { + type<T>(); + + return detail::meta_kernel_literal<T>(value); + } + + template<class T> + detail::meta_kernel_variable<T> expr(const std::string &expr) const + { + type<T>(); + + return detail::meta_kernel_variable<T>(expr); + } + + // define stream operators for scalar and vector types + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(char) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(uchar) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(short) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(ushort) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(int) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(uint) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(long) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(ulong) + BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(double) + + // define stream operators for float scalar and vector types + meta_kernel& operator<<(const float &x) + { + m_source << std::showpoint << x << 'f'; + return *this; + } + + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float2_) + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float4_) + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float8_) + BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float16_) + + // define stream operators for variable types + template<class T> + meta_kernel& operator<<(const meta_kernel_variable<T> &variable) + { + return *this << variable.name(); + } + + // define stream operators for literal types + template<class T> + meta_kernel& operator<<(const meta_kernel_literal<T> &literal) + { + return *this << literal.value(); + } + + meta_kernel& operator<<(const meta_kernel_literal<bool> &literal) + { + return *this << (literal.value() ? "true" : "false"); + } + + meta_kernel& operator<<(const meta_kernel_literal<char> &literal) + { + const char c = literal.value(); + + switch(c){ + // control characters + case '\0': + return *this << "'\\0'"; + case '\a': + return *this << "'\\a'"; + case '\b': + return *this << "'\\b'"; + case '\t': + return *this << "'\\t'"; + case '\n': + return *this << "'\\n'"; + case '\v': + return *this << "'\\v'"; + case '\f': + return *this << "'\\f'"; + case '\r': + return *this << "'\\r'"; + + // characters which need escaping + case '\"': + case '\'': + case '\?': + case '\\': + return *this << "'\\" << c << "'"; + + // all other characters + default: + return *this << "'" << c << "'"; + } + } + + meta_kernel& operator<<(const meta_kernel_literal<signed char> &literal) + { + return *this << lit<char>(literal.value()); + } + + meta_kernel& operator<<(const meta_kernel_literal<unsigned char> &literal) + { + return *this << uint_(literal.value()); + } + + // define stream operators for strings + meta_kernel& operator<<(char ch) + { + m_source << ch; + return *this; + } + + meta_kernel& operator<<(const char *string) + { + m_source << string; + return *this; + } + + meta_kernel& operator<<(const std::string &string) + { + m_source << string; + return *this; + } + + template<class T> + static detail::meta_kernel_variable<T> make_var(const std::string &name) + { + return detail::meta_kernel_variable<T>(name); + } + + template<class T> + static detail::meta_kernel_literal<T> make_lit(const T &value) + { + return detail::meta_kernel_literal<T>(value); + } + + template<class T> + static detail::meta_kernel_variable<T> make_expr(const std::string &expr) + { + return detail::meta_kernel_variable<T>(expr); + } + + event exec(command_queue &queue) + { + return exec_1d(queue, 0, 1); + } + + event exec_1d(command_queue &queue, + size_t global_work_offset, + size_t global_work_size) + { + const context &context = queue.get_context(); + + ::boost::compute::kernel kernel = compile(context); + + return queue.enqueue_1d_range_kernel( + kernel, + global_work_offset, + global_work_size, + 0 + ); + } + + event exec_1d(command_queue &queue, + size_t global_work_offset, + size_t global_work_size, + size_t local_work_size) + { + const context &context = queue.get_context(); + + ::boost::compute::kernel kernel = compile(context); + + return queue.enqueue_1d_range_kernel( + kernel, + global_work_offset, + global_work_size, + local_work_size + ); + } + + template<class T> + std::string get_buffer_identifier(const buffer &buffer, + const memory_object::address_space address_space = + memory_object::global_memory) + { + // check if we've already seen buffer + for(size_t i = 0; i < m_stored_buffers.size(); i++){ + const detail::meta_kernel_buffer_info &bi = m_stored_buffers[i]; + + if(bi.m_mem == buffer.get() && + bi.address_space == address_space){ + return bi.identifier; + } + } + + // create a new binding + std::string identifier = + "_buf" + lexical_cast<std::string>(m_stored_buffers.size()); + size_t index = add_arg<T *>(address_space, identifier); + + // store new buffer info + m_stored_buffers.push_back( + detail::meta_kernel_buffer_info(buffer, identifier, address_space, index)); + + return identifier; + } + + std::string get_image_identifier(const char *qualifiers, const image2d &image) + { + size_t index = add_arg_with_qualifiers<image2d>(qualifiers, "image"); + + set_arg(index, image); + + return "image"; + } + + std::string get_sampler_identifier(bool normalized_coords, + cl_addressing_mode addressing_mode, + cl_filter_mode filter_mode) + { + (void) normalized_coords; + (void) addressing_mode; + (void) filter_mode; + + m_pragmas += "const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |\n" + " CLK_ADDRESS_NONE |\n" + " CLK_FILTER_NEAREST;\n"; + + return "sampler"; + } + + template<class Expr> + static std::string expr_to_string(const Expr &expr) + { + meta_kernel tmp((std::string())); + tmp << expr; + return tmp.m_source.str(); + } + + template<class Predicate> + detail::invoked_function<bool, boost::tuple<Predicate> > if_(Predicate pred) const + { + return detail::invoked_function<bool, boost::tuple<Predicate> >( + "if", std::string(), boost::make_tuple(pred) + ); + } + + template<class Predicate> + detail::invoked_function<bool, boost::tuple<Predicate> > else_if_(Predicate pred) const + { + return detail::invoked_function<bool, boost::tuple<Predicate> >( + "else if", std::string(), boost::make_tuple(pred) + ); + } + + detail::meta_kernel_variable<cl_uint> get_global_id(size_t dim) const + { + return expr<cl_uint>("get_global_id(" + lexical_cast<std::string>(dim) + ")"); + } + + void add_function(const std::string &name, const std::string &source) + { + if(m_external_function_names.count(name)){ + return; + } + + m_external_function_names.insert(name); + m_external_function_source << source << "\n"; + } + + void add_function(const std::string &name, + const std::string &source, + const std::map<std::string, std::string> &definitions) + { + typedef std::map<std::string, std::string>::const_iterator iter; + + std::stringstream s; + + // add #define's + for(iter i = definitions.begin(); i != definitions.end(); i++){ + s << "#define " << i->first; + if(!i->second.empty()){ + s << " " << i->second; + } + s << "\n"; + } + + s << source << "\n"; + + // add #undef's + for(iter i = definitions.begin(); i != definitions.end(); i++){ + s << "#undef " << i->first << "\n"; + } + + add_function(name, s.str()); + } + + template<class Type> + void add_type_declaration(const std::string &declaration) + { + const char *name = type_name<Type>(); + + // check if the type has already been declared + std::string source = m_type_declaration_source.str(); + if(source.find(name) != std::string::npos){ + return; + } + + m_type_declaration_source << declaration; + } + + template<class Type> + void inject_type() const + { + inject_type_impl<Type>()(const_cast<meta_kernel &>(*this)); + } + + // the insert_function_call() method inserts a call to a function with + // the given name tuple of argument values. + template<class ArgTuple> + void insert_function_call(const std::string &name, const ArgTuple &args) + { + *this << name << '('; + insert_function_call_args(args); + *this << ')'; + } + + // the insert_function_call_args() method takes a tuple of argument values + // and inserts them into the source string with a comma in-between each. + // this is useful for creating function calls given a tuple of values. + void insert_function_call_args(const boost::tuple<>&) + { + } + + #define BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE(z, n, unused) \ + inject_type<BOOST_PP_CAT(T, n)>(); + + #define BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG(z, n, unused) \ + << boost::get<BOOST_PP_DEC(n)>(args) << ", " + + #define BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS(z, n, unused) \ + template<BOOST_PP_ENUM_PARAMS(n, class T)> \ + void insert_function_call_args( \ + const boost::tuple<BOOST_PP_ENUM_PARAMS(n, T)> &args \ + ) \ + { \ + BOOST_PP_REPEAT_FROM_TO( \ + 0, n, BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE, ~ \ + ) \ + *this \ + BOOST_PP_REPEAT_FROM_TO( \ + 1, n, BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG, ~ \ + ) \ + << boost::get<BOOST_PP_DEC(n)>(args); \ + } + + BOOST_PP_REPEAT_FROM_TO( + 1, BOOST_COMPUTE_MAX_ARITY, BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS, ~ + ) + + #undef BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE + #undef BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG + #undef BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS + + static const char* address_space_prefix(const memory_object::address_space value) + { + switch(value){ + case memory_object::global_memory: return "__global"; + case memory_object::local_memory: return "__local"; + case memory_object::private_memory: return "__private"; + case memory_object::constant_memory: return "__constant"; + }; + + return 0; // unreachable + } + +private: + template<class T> + size_t add_arg_with_qualifiers(const char *qualifiers, const std::string &name) + { + size_t index = add_arg<T>(name); + + // update argument type declaration with qualifiers + std::stringstream s; + s << qualifiers << " " << m_args[index]; + m_args[index] = s.str(); + + return index; + } + +private: + std::string m_name; + std::stringstream m_source; + std::stringstream m_external_function_source; + std::stringstream m_type_declaration_source; + std::set<std::string> m_external_function_names; + std::vector<std::string> m_args; + std::string m_pragmas; + std::vector<detail::meta_kernel_stored_arg> m_stored_args; + std::vector<detail::meta_kernel_buffer_info> m_stored_buffers; +}; + +template<class ResultType, class ArgTuple> +inline meta_kernel& +operator<<(meta_kernel &kernel, const invoked_function<ResultType, ArgTuple> &expr) +{ + if(!expr.source().empty()){ + kernel.add_function(expr.name(), expr.source(), expr.definitions()); + } + + kernel.insert_function_call(expr.name(), expr.args()); + + return kernel; +} + +template<class ResultType, class ArgTuple, class CaptureTuple> +inline meta_kernel& +operator<<(meta_kernel &kernel, + const invoked_closure<ResultType, ArgTuple, CaptureTuple> &expr) +{ + if(!expr.source().empty()){ + kernel.add_function(expr.name(), expr.source(), expr.definitions()); + } + + kernel << expr.name() << '('; + kernel.insert_function_call_args(expr.args()); + kernel << ", "; + kernel.insert_function_call_args(expr.capture()); + kernel << ')'; + + return kernel; +} + +template<class Arg1, class Arg2, class Result> +inline meta_kernel& operator<<(meta_kernel &kernel, + const invoked_binary_operator<Arg1, + Arg2, + Result> &expr) +{ + return kernel << "((" << expr.arg1() << ")" + << expr.op() + << "(" << expr.arg2() << "))"; +} + +template<class T, class IndexExpr> +inline meta_kernel& operator<<(meta_kernel &kernel, + const detail::device_ptr_index_expr<T, IndexExpr> &expr) +{ + if(expr.m_index == 0){ + return kernel << + kernel.get_buffer_identifier<T>(expr.m_buffer) << + '[' << expr.m_expr << ']'; + } + else { + return kernel << + kernel.get_buffer_identifier<T>(expr.m_buffer) << + '[' << expr.m_index << "+(" << expr.m_expr << ")]"; + } +} + +template<class T1, class T2, class IndexExpr> +inline meta_kernel& operator<<(meta_kernel &kernel, + const detail::device_ptr_index_expr<std::pair<T1, T2>, IndexExpr> &expr) +{ + typedef std::pair<T1, T2> T; + + if(expr.m_index == 0){ + return kernel << + kernel.get_buffer_identifier<T>(expr.m_buffer) << + '[' << expr.m_expr << ']'; + } + else { + return kernel << + kernel.get_buffer_identifier<T>(expr.m_buffer) << + '[' << expr.m_index << "+(" << expr.m_expr << ")]"; + } +} + +template<class Predicate, class Arg> +inline meta_kernel& operator<<(meta_kernel &kernel, + const invoked_unary_negate_function<Predicate, + Arg> &expr) +{ + return kernel << "!(" << expr.pred()(expr.expr()) << ')'; +} + +template<class Predicate, class Arg1, class Arg2> +inline meta_kernel& operator<<(meta_kernel &kernel, + const invoked_binary_negate_function<Predicate, + Arg1, + Arg2> &expr) +{ + return kernel << "!(" << expr.pred()(expr.expr1(), expr.expr2()) << ')'; +} + +// get<N>() for vector types +template<size_t N, class Arg, class T> +inline meta_kernel& operator<<(meta_kernel &kernel, + const invoked_get<N, Arg, T> &expr) +{ + BOOST_STATIC_ASSERT(N < 16); + + if(N < 10){ + return kernel << expr.m_arg << ".s" << uint_(N); + } + else if(N < 16){ +#ifdef _MSC_VER +# pragma warning(push) +# pragma warning(disable: 4307) +#endif + return kernel << expr.m_arg << ".s" << char('a' + (N - 10)); +#ifdef _MSC_VER +# pragma warning(pop) +#endif + } + + return kernel; +} + +template<class T, class Arg> +inline meta_kernel& operator<<(meta_kernel &kernel, + const invoked_field<T, Arg> &expr) +{ + return kernel << expr.m_arg << "." << expr.m_field; +} + +template<class T, class Arg> +inline meta_kernel& operator<<(meta_kernel &k, + const invoked_as<T, Arg> &expr) +{ + return k << "as_" << type_name<T>() << "(" << expr.m_arg << ")"; +} + +template<class T, class Arg> +inline meta_kernel& operator<<(meta_kernel &k, + const invoked_convert<T, Arg> &expr) +{ + return k << "convert_" << type_name<T>() << "(" << expr.m_arg << ")"; +} + +template<class T, class Arg> +inline meta_kernel& operator<<(meta_kernel &k, + const invoked_identity<T, Arg> &expr) +{ + return k << expr.m_arg; +} + +template<> +struct inject_type_impl<double_> +{ + void operator()(meta_kernel &kernel) + { + kernel.add_extension_pragma("cl_khr_fp64", "enable"); + } +}; + +template<class Scalar, size_t N> +struct inject_type_impl<vector_type<Scalar, N> > +{ + void operator()(meta_kernel &kernel) + { + kernel.inject_type<Scalar>(); + } +}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_META_KERNEL_HPP diff --git a/boost/compute/detail/mpl_vector_to_tuple.hpp b/boost/compute/detail/mpl_vector_to_tuple.hpp new file mode 100644 index 0000000000..292a6e36e1 --- /dev/null +++ b/boost/compute/detail/mpl_vector_to_tuple.hpp @@ -0,0 +1,65 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_MPL_VECTOR_TO_TUPLE_HPP +#define BOOST_COMPUTE_DETAIL_MPL_VECTOR_TO_TUPLE_HPP + +#include <boost/mpl/copy.hpp> +#include <boost/mpl/vector.hpp> +#include <boost/tuple/tuple.hpp> +#include <boost/fusion/include/mpl.hpp> +#include <boost/fusion/adapted/boost_tuple.hpp> +#include <boost/preprocessor/repetition.hpp> + +#include <boost/compute/config.hpp> + +namespace boost { +namespace compute { +namespace detail { + +namespace mpl = boost::mpl; + +template<class Vector, size_t N> +struct mpl_vector_to_tuple_impl; + +#define BOOST_COMPUTE_PRINT_ELEM(z, n, unused) \ + typename mpl::at_c<Vector, n>::type + +#define BOOST_COMPUTE_VEC2TUP(z, n, unused) \ +template<class Vector> \ +struct mpl_vector_to_tuple_impl<Vector, n> \ +{ \ + typedef typename \ + boost::tuple< \ + BOOST_PP_ENUM(n, BOOST_COMPUTE_PRINT_ELEM, ~) \ + > type; \ +}; + +BOOST_PP_REPEAT_FROM_TO(1, BOOST_COMPUTE_MAX_ARITY, BOOST_COMPUTE_VEC2TUP, ~) + +#undef BOOST_COMPUTE_VEC2TUP +#undef BOOST_COMPUTE_PRINT_ELEM + +// meta-function which converts a mpl::vector to a boost::tuple +template<class Vector> +struct mpl_vector_to_tuple +{ + typedef typename + mpl_vector_to_tuple_impl< + Vector, + mpl::size<Vector>::value + >::type type; +}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_MPL_VECTOR_TO_TUPLE_HPP diff --git a/boost/compute/detail/nvidia_compute_capability.hpp b/boost/compute/detail/nvidia_compute_capability.hpp new file mode 100644 index 0000000000..3f859562bd --- /dev/null +++ b/boost/compute/detail/nvidia_compute_capability.hpp @@ -0,0 +1,60 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_NVIDIA_COMPUTE_CAPABILITY_HPP +#define BOOST_COMPUTE_DETAIL_NVIDIA_COMPUTE_CAPABILITY_HPP + +#include <boost/compute/device.hpp> + +#ifdef BOOST_COMPUTE_HAVE_HDR_CL_EXT + #include <CL/cl_ext.h> +#endif + +namespace boost { +namespace compute { +namespace detail { + +#ifdef CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV + #define BOOST_COMPUTE_CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV +#else + #define BOOST_COMPUTE_CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV 0x4000 +#endif + +#ifdef CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV + #define BOOST_COMPUTE_CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV +#else + #define BOOST_COMPUTE_CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV 0x4001 +#endif + +inline void get_nvidia_compute_capability(const device &device, int &major, int &minor) +{ + if(!device.supports_extension("cl_nv_device_attribute_query")){ + major = minor = 0; + return; + } + + major = device.get_info<uint_>(BOOST_COMPUTE_CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV); + minor = device.get_info<uint_>(BOOST_COMPUTE_CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV); +} + +inline bool check_nvidia_compute_capability(const device &device, int major, int minor) +{ + int actual_major, actual_minor; + get_nvidia_compute_capability(device, actual_major, actual_minor); + + return actual_major > major || + (actual_major == major && actual_minor >= minor); +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_NVIDIA_COMPUTE_CAPABILITY_HPP diff --git a/boost/compute/detail/parameter_cache.hpp b/boost/compute/detail/parameter_cache.hpp new file mode 100644 index 0000000000..2a856311e1 --- /dev/null +++ b/boost/compute/detail/parameter_cache.hpp @@ -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://boostorg.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 ¶meter, 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 ¶meter, 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 diff --git a/boost/compute/detail/path.hpp b/boost/compute/detail/path.hpp new file mode 100644 index 0000000000..ec8760eaf9 --- /dev/null +++ b/boost/compute/detail/path.hpp @@ -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://boostorg.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 diff --git a/boost/compute/detail/print_range.hpp b/boost/compute/detail/print_range.hpp new file mode 100644 index 0000000000..bfe02f6828 --- /dev/null +++ b/boost/compute/detail/print_range.hpp @@ -0,0 +1,82 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_PRINT_RANGE_HPP +#define BOOST_COMPUTE_DETAIL_PRINT_RANGE_HPP + +#include <vector> +#include <iostream> +#include <iterator> + +#include <boost/compute/algorithm/copy.hpp> +#include <boost/compute/container/vector.hpp> +#include <boost/compute/detail/is_buffer_iterator.hpp> +#include <boost/compute/detail/iterator_range_size.hpp> + +namespace boost { +namespace compute { +namespace detail { + +template<class InputIterator> +inline void print_range(InputIterator first, + InputIterator last, + command_queue &queue, + typename boost::enable_if< + is_buffer_iterator<InputIterator> + >::type* = 0) +{ + typedef typename + std::iterator_traits<InputIterator>::value_type + value_type; + + const size_t size = iterator_range_size(first, last); + + // copy values to temporary vector on the host + std::vector<value_type> tmp(size); + ::boost::compute::copy(first, last, tmp.begin(), queue); + + // print values + std::cout << "[ "; + for(size_t i = 0; i < size; i++){ + std::cout << tmp[i]; + if(i != size - 1){ + std::cout << ", "; + } + } + std::cout << " ]" << std::endl; +} + +template<class InputIterator> +inline void print_range(InputIterator first, + InputIterator last, + command_queue &queue, + typename boost::enable_if_c< + !is_buffer_iterator<InputIterator>::value + >::type* = 0) +{ + typedef typename + std::iterator_traits<InputIterator>::value_type + value_type; + + const context &context = queue.get_context(); + const size_t size = iterator_range_size(first, last); + + // copy values to temporary vector on the device + ::boost::compute::vector<value_type> tmp(size, context); + ::boost::compute::copy(first, last, tmp.begin(), queue); + + print_range(tmp.begin(), tmp.end(), queue); +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_PRINT_RANGE_HPP diff --git a/boost/compute/detail/read_write_single_value.hpp b/boost/compute/detail/read_write_single_value.hpp new file mode 100644 index 0000000000..fde40d946c --- /dev/null +++ b/boost/compute/detail/read_write_single_value.hpp @@ -0,0 +1,77 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_READ_WRITE_SINGLE_VALUE_HPP +#define BOOST_COMPUTE_DETAIL_READ_WRITE_SINGLE_VALUE_HPP + +#include <boost/throw_exception.hpp> + +#include <boost/compute/buffer.hpp> +#include <boost/compute/exception.hpp> +#include <boost/compute/command_queue.hpp> + +namespace boost { +namespace compute { +namespace detail { + +// reads and returns a single value at index in the buffer +template<class T> +inline T read_single_value(const buffer &buffer, + size_t index, + command_queue &queue) +{ + BOOST_ASSERT(index < buffer.size() / sizeof(T)); + BOOST_ASSERT(buffer.get_context() == queue.get_context()); + + T value; + queue.enqueue_read_buffer(buffer, + sizeof(T) * index, + sizeof(T), + &value); + return value; +} + +// reads and returns a the first value in the buffer +template<class T> +inline T read_single_value(const buffer &buffer, command_queue &queue) +{ + return read_single_value<T>(buffer, 0, queue); +} + +// writes a single value at index to the buffer +template<class T> +inline void write_single_value(const T &value, + const buffer &buffer, + size_t index, + command_queue &queue) +{ + BOOST_ASSERT(index < buffer.size() / sizeof(T)); + BOOST_ASSERT(buffer.get_context() == queue.get_context()); + + queue.enqueue_write_buffer(buffer, + index * sizeof(T), + sizeof(T), + &value); +} + +// writes value to the first location in buffer +template<class T> +inline void write_single_value(const T &value, + const buffer &buffer, + command_queue &queue) +{ + write_single_value<T>(value, buffer, 0, queue); +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_READ_WRITE_SINGLE_VALUE_HPP diff --git a/boost/compute/detail/sha1.hpp b/boost/compute/detail/sha1.hpp new file mode 100644 index 0000000000..5685fa4407 --- /dev/null +++ b/boost/compute/detail/sha1.hpp @@ -0,0 +1,53 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_SHA1_HPP +#define BOOST_COMPUTE_DETAIL_SHA1_HPP + +#include <sstream> +#include <iomanip> +#include <boost/uuid/sha1.hpp> + +namespace boost { +namespace compute { +namespace detail { + +// Accumulates SHA1 hash of the passed strings. +class sha1 { + public: + sha1(const std::string &s = "") { + if (!s.empty()) this->process(s); + } + + sha1& process(const std::string &s) { + h.process_bytes(s.c_str(), s.size()); + return *this; + } + + operator std::string() { + unsigned int digest[5]; + h.get_digest(digest); + + std::ostringstream buf; + for(int i = 0; i < 5; ++i) + buf << std::hex << std::setfill('0') << std::setw(8) << digest[i]; + + return buf.str(); + } + private: + boost::uuids::detail::sha1 h; +}; + +} // end detail namespace +} // end compute namespace +} // end boost namespace + + +#endif // BOOST_COMPUTE_DETAIL_SHA1_HPP diff --git a/boost/compute/detail/variadic_macros.hpp b/boost/compute/detail/variadic_macros.hpp new file mode 100644 index 0000000000..60f44bd1a8 --- /dev/null +++ b/boost/compute/detail/variadic_macros.hpp @@ -0,0 +1,35 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_VARIDAIC_MACROS_HPP +#define BOOST_COMPUTE_DETAIL_VARIDAIC_MACROS_HPP + +#include <boost/preprocessor/cat.hpp> +#include <boost/preprocessor/config/config.hpp> +#include <boost/preprocessor/tuple/to_seq.hpp> + +#if BOOST_PP_VARIADICS == 1 +# include <boost/preprocessor/variadic/size.hpp> +#endif + +#ifdef BOOST_PP_VARIADIC_SIZE +# define BOOST_COMPUTE_PP_VARIADIC_SIZE BOOST_PP_VARIADIC_SIZE +#else +# define BOOST_COMPUTE_PP_VARIADIC_SIZE(...) BOOST_COMPUTE_PP_VARIADIC_SIZE_I(__VA_ARGS__, 64, 63, 62, 61, 60, 59, 58, 57, 56, 55, 54, 53, 52, 51, 50, 49, 48, 47, 46, 45, 44, 43, 42, 41, 40, 39, 38, 37, 36, 35, 34, 33, 32, 31, 30, 29, 28, 27, 26, 25, 24, 23, 22, 21, 20, 19, 18, 17, 16, 15, 14, 13, 12, 11, 10, 9, 8, 7, 6, 5, 4, 3, 2, 1,) +# define BOOST_COMPUTE_PP_VARIADIC_SIZE_I(e0, e1, e2, e3, e4, e5, e6, e7, e8, e9, e10, e11, e12, e13, e14, e15, e16, e17, e18, e19, e20, e21, e22, e23, e24, e25, e26, e27, e28, e29, e30, e31, e32, e33, e34, e35, e36, e37, e38, e39, e40, e41, e42, e43, e44, e45, e46, e47, e48, e49, e50, e51, e52, e53, e54, e55, e56, e57, e58, e59, e60, e61, e62, e63, size, ...) size +#endif + +#define BOOST_COMPUTE_PP_TUPLE_SIZE(tuple) \ + BOOST_COMPUTE_PP_VARIADIC_SIZE tuple + +#define BOOST_COMPUTE_PP_TUPLE_TO_SEQ(tuple) \ + BOOST_PP_TUPLE_TO_SEQ(BOOST_COMPUTE_PP_TUPLE_SIZE(tuple), tuple) + +#endif // BOOST_COMPUTE_DETAIL_VARIDAIC_MACROS_HPP diff --git a/boost/compute/detail/vendor.hpp b/boost/compute/detail/vendor.hpp new file mode 100644 index 0000000000..0aa9c9c0d4 --- /dev/null +++ b/boost/compute/detail/vendor.hpp @@ -0,0 +1,38 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_VENDOR_HPP +#define BOOST_COMPUTE_DETAIL_VENDOR_HPP + +#include <boost/compute/device.hpp> +#include <boost/compute/platform.hpp> + +namespace boost { +namespace compute { +namespace detail { + +// returns true if the device is an nvidia gpu +inline bool is_nvidia_device(const device &device) +{ + std::string nvidia("NVIDIA"); + return device.vendor().compare(0, nvidia.size(), nvidia) == 0; +} + +// returns true if the device is an amd cpu or gpu +inline bool is_amd_device(const device &device) +{ + return device.platform().vendor() == "Advanced Micro Devices, Inc."; +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_VENDOR_HPP diff --git a/boost/compute/detail/work_size.hpp b/boost/compute/detail/work_size.hpp new file mode 100644 index 0000000000..552d797b8b --- /dev/null +++ b/boost/compute/detail/work_size.hpp @@ -0,0 +1,37 @@ +//---------------------------------------------------------------------------// +// 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://boostorg.github.com/compute for more information. +//---------------------------------------------------------------------------// + +#ifndef BOOST_COMPUTE_DETAIL_WORK_SIZE_HPP +#define BOOST_COMPUTE_DETAIL_WORK_SIZE_HPP + +#include <cmath> + +namespace boost { +namespace compute { +namespace detail { + +// Given a total number of values (count), a number of values to +// process per thread (vtp), and a number of threads to execute per +// block (tpb), this function returns the global work size to be +// passed to clEnqueueNDRangeKernel() for a 1D algorithm. +inline size_t calculate_work_size(size_t count, size_t vpt, size_t tpb) +{ + size_t work_size = static_cast<size_t>(std::ceil(float(count) / vpt)); + if(work_size % tpb != 0){ + work_size += tpb - work_size % tpb; + } + return work_size; +} + +} // end detail namespace +} // end compute namespace +} // end boost namespace + +#endif // BOOST_COMPUTE_DETAIL_WORK_SIZE_HPP |