Commit 3f0a1191 authored by Dominik Charousset's avatar Dominik Charousset

used limited vector for OpenCL dimensions

OpenCL interface now uses vectors with a maximum size of 3, because
the dimensions settings of OpenCL assume a one, two, or three dimensional space
parent e2832362
...@@ -565,7 +565,7 @@ WARN_LOGFILE = ...@@ -565,7 +565,7 @@ WARN_LOGFILE =
# directories like "/usr/src/myproject". Separate the files or directories # directories like "/usr/src/myproject". Separate the files or directories
# with spaces. # with spaces.
INPUT = @CMAKE_HOME_DIRECTORY@/cppa/ @CMAKE_HOME_DIRECTORY@/cppa/util @CMAKE_HOME_DIRECTORY@/cppa/intrusive @CMAKE_HOME_DIRECTORY@/cppa/network INPUT = @CMAKE_HOME_DIRECTORY@/cppa/ @CMAKE_HOME_DIRECTORY@/cppa/util @CMAKE_HOME_DIRECTORY@/cppa/intrusive @CMAKE_HOME_DIRECTORY@/cppa/network @CMAKE_HOME_DIRECTORY@/cppa/opencl
# This tag can be used to specify the character encoding of the source files # This tag can be used to specify the character encoding of the source files
# that doxygen parses. Internally doxygen uses the UTF-8 encoding, which is # that doxygen parses. Internally doxygen uses the UTF-8 encoding, which is
......
...@@ -156,7 +156,6 @@ cppa/util/dptr.hpp ...@@ -156,7 +156,6 @@ cppa/util/dptr.hpp
cppa/util/duration.hpp cppa/util/duration.hpp
cppa/util/element_at.hpp cppa/util/element_at.hpp
cppa/util/fiber.hpp cppa/util/fiber.hpp
cppa/util/fixed_vector.hpp
cppa/util/get_result_type.hpp cppa/util/get_result_type.hpp
cppa/util/int_list.hpp cppa/util/int_list.hpp
cppa/util/is_array_of.hpp cppa/util/is_array_of.hpp
...@@ -170,6 +169,7 @@ cppa/util/is_manipulator.hpp ...@@ -170,6 +169,7 @@ cppa/util/is_manipulator.hpp
cppa/util/is_mutable_ref.hpp cppa/util/is_mutable_ref.hpp
cppa/util/is_primitive.hpp cppa/util/is_primitive.hpp
cppa/util/left_or_right.hpp cppa/util/left_or_right.hpp
cppa/util/limited_vector.hpp
cppa/util/producer_consumer_list.hpp cppa/util/producer_consumer_list.hpp
cppa/util/pt_dispatch.hpp cppa/util/pt_dispatch.hpp
cppa/util/pt_token.hpp cppa/util/pt_token.hpp
......
...@@ -42,7 +42,7 @@ ...@@ -42,7 +42,7 @@
#include "cppa/ref_counted.hpp" #include "cppa/ref_counted.hpp"
#include "cppa/util/at.hpp" #include "cppa/util/at.hpp"
#include "cppa/util/fixed_vector.hpp" #include "cppa/util/limited_vector.hpp"
#include "cppa/util/is_comparable.hpp" #include "cppa/util/is_comparable.hpp"
#include "cppa/util/compare_tuples.hpp" #include "cppa/util/compare_tuples.hpp"
#include "cppa/util/is_legal_tuple_type.hpp" #include "cppa/util/is_legal_tuple_type.hpp"
...@@ -155,7 +155,7 @@ class cow_tuple<Head, Tail...> { ...@@ -155,7 +155,7 @@ class cow_tuple<Head, Tail...> {
} }
static cow_tuple from(cow_ptr_type ptr, static cow_tuple from(cow_ptr_type ptr,
const util::fixed_vector<size_t, num_elements>& mv) { const util::limited_vector<size_t, num_elements>& mv) {
return {priv_ctor{}, decorated_type::create(std::move(ptr), mv)}; return {priv_ctor{}, decorated_type::create(std::move(ptr), mv)};
} }
......
...@@ -141,6 +141,9 @@ ...@@ -141,6 +141,9 @@
* @namespace cppa::intrusive * @namespace cppa::intrusive
* @brief Contains intrusive container implementations. * @brief Contains intrusive container implementations.
* *
* @namespace cppa::opencl
* @brief Contains all classes of libcppa's OpenCL binding (optional).
*
* @namespace cppa::network * @namespace cppa::network
* @brief Contains all network related classes. * @brief Contains all network related classes.
* *
......
...@@ -40,7 +40,7 @@ ...@@ -40,7 +40,7 @@
#include "cppa/uniform_type_info.hpp" #include "cppa/uniform_type_info.hpp"
#include "cppa/util/type_list.hpp" #include "cppa/util/type_list.hpp"
#include "cppa/util/fixed_vector.hpp" #include "cppa/util/limited_vector.hpp"
#include "cppa/detail/tuple_vals.hpp" #include "cppa/detail/tuple_vals.hpp"
#include "cppa/detail/abstract_tuple.hpp" #include "cppa/detail/abstract_tuple.hpp"
...@@ -58,7 +58,7 @@ class decorated_tuple : public abstract_tuple { ...@@ -58,7 +58,7 @@ class decorated_tuple : public abstract_tuple {
public: public:
typedef util::fixed_vector<size_t, sizeof...(Ts)> vector_type; typedef util::limited_vector<size_t, sizeof...(Ts)> vector_type;
typedef cow_ptr<abstract_tuple> cow_pointer_type; typedef cow_ptr<abstract_tuple> cow_pointer_type;
......
...@@ -61,7 +61,7 @@ struct matcher<wildcard_position::nil, Tuple, T...> { ...@@ -61,7 +61,7 @@ struct matcher<wildcard_position::nil, Tuple, T...> {
} }
static inline bool tmatch(const Tuple& tup, static inline bool tmatch(const Tuple& tup,
util::fixed_vector<size_t, sizeof...(T)>& mv) { util::limited_vector<size_t, sizeof...(T)>& mv) {
if (tmatch(tup)) { if (tmatch(tup)) {
mv.resize(sizeof...(T)); mv.resize(sizeof...(T));
std::iota(mv.begin(), mv.end(), 0); std::iota(mv.begin(), mv.end(), 0);
...@@ -86,7 +86,7 @@ struct matcher<wildcard_position::trailing, Tuple, T...> { ...@@ -86,7 +86,7 @@ struct matcher<wildcard_position::trailing, Tuple, T...> {
} }
static inline bool tmatch(const Tuple& tup, static inline bool tmatch(const Tuple& tup,
util::fixed_vector<size_t, size>& mv) { util::limited_vector<size_t, size>& mv) {
if (tmatch(tup)) { if (tmatch(tup)) {
mv.resize(size); mv.resize(size);
std::iota(mv.begin(), mv.end(), 0); std::iota(mv.begin(), mv.end(), 0);
...@@ -101,7 +101,7 @@ struct matcher<wildcard_position::leading, Tuple, anything> { ...@@ -101,7 +101,7 @@ struct matcher<wildcard_position::leading, Tuple, anything> {
static inline bool tmatch(const Tuple&) { static inline bool tmatch(const Tuple&) {
return true; return true;
} }
static inline bool tmatch(const Tuple&, util::fixed_vector<size_t, 0>&) { static inline bool tmatch(const Tuple&, util::limited_vector<size_t, 0>&) {
return true; return true;
} }
}; };
...@@ -123,7 +123,7 @@ struct matcher<wildcard_position::leading, Tuple, T...> { ...@@ -123,7 +123,7 @@ struct matcher<wildcard_position::leading, Tuple, T...> {
} }
static inline bool tmatch(const Tuple& tup, static inline bool tmatch(const Tuple& tup,
util::fixed_vector<size_t, size>& mv) { util::limited_vector<size_t, size>& mv) {
if (tmatch(tup)) { if (tmatch(tup)) {
mv.resize(size); mv.resize(size);
std::iota(mv.begin(), mv.end(), tup.size() - size); std::iota(mv.begin(), mv.end(), tup.size() - size);
...@@ -164,7 +164,7 @@ struct matcher<wildcard_position::in_between, Tuple, T...> { ...@@ -164,7 +164,7 @@ struct matcher<wildcard_position::in_between, Tuple, T...> {
} }
static inline bool tmatch(const Tuple& tup, static inline bool tmatch(const Tuple& tup,
util::fixed_vector<size_t, size - 1>& mv) { util::limited_vector<size_t, size - 1>& mv) {
if (tmatch(tup)) { if (tmatch(tup)) {
// first range // first range
mv.resize(size - 1); mv.resize(size - 1);
...@@ -263,7 +263,7 @@ struct match_impl { ...@@ -263,7 +263,7 @@ struct match_impl {
template<size_t Size> template<size_t Size>
static inline bool _(const Tuple& tup, static inline bool _(const Tuple& tup,
util::fixed_vector<size_t, Size>& mv) { util::limited_vector<size_t, Size>& mv) {
return matcher<PC, Tuple, Ts...>::tmatch(tup, mv); return matcher<PC, Tuple, Ts...>::tmatch(tup, mv);
} }
}; };
...@@ -279,7 +279,7 @@ struct match_impl<wildcard_position::multiple, Tuple, Ts...> { ...@@ -279,7 +279,7 @@ struct match_impl<wildcard_position::multiple, Tuple, Ts...> {
template<size_t Size> template<size_t Size>
static inline bool _(const Tuple& tup, static inline bool _(const Tuple& tup,
util::fixed_vector<size_t, Size>& mv) { util::limited_vector<size_t, Size>& mv) {
return matcher<PC, Tuple, Ts...>::tmatch(tup, mv); return matcher<PC, Tuple, Ts...>::tmatch(tup, mv);
} }
...@@ -311,7 +311,7 @@ bool matches(const any_tuple& tup) { ...@@ -311,7 +311,7 @@ bool matches(const any_tuple& tup) {
*/ */
template<typename... Ts> template<typename... Ts>
bool matches(const any_tuple& tup, bool matches(const any_tuple& tup,
util::fixed_vector< util::limited_vector<
size_t, size_t,
util::tl_count_not< util::tl_count_not<
util::type_list<Ts...>, util::type_list<Ts...>,
...@@ -329,7 +329,7 @@ inline bool matches(const any_tuple& tup, const util::type_list<Ts...>&) { ...@@ -329,7 +329,7 @@ inline bool matches(const any_tuple& tup, const util::type_list<Ts...>&) {
template<typename... Ts> template<typename... Ts>
inline bool matches(const any_tuple& tup, const util::type_list<Ts...>&, inline bool matches(const any_tuple& tup, const util::type_list<Ts...>&,
util::fixed_vector< util::limited_vector<
size_t, size_t,
util::tl_count_not< util::tl_count_not<
util::type_list<Ts...>, util::type_list<Ts...>,
......
...@@ -57,7 +57,7 @@ struct tuple_cast_impl { ...@@ -57,7 +57,7 @@ struct tuple_cast_impl {
static constexpr size_t first_wc = static constexpr size_t first_wc =
static_cast<size_t>( static_cast<size_t>(
util::tl_find<util::type_list<T...>, anything>::value); util::tl_find<util::type_list<T...>, anything>::value);
typedef util::fixed_vector<size_t, size> mapping_vector; typedef util::limited_vector<size_t, size> mapping_vector;
static inline option<Result> safe(any_tuple& tup) { static inline option<Result> safe(any_tuple& tup) {
mapping_vector mv; mapping_vector mv;
if (matches<T...>(tup, mv)) return {Result::from(std::move(tup.vals()), if (matches<T...>(tup, mv)) return {Result::from(std::move(tup.vals()),
......
...@@ -98,7 +98,7 @@ struct invoke_policy_impl : invoke_policy_base<FilteredPattern> { ...@@ -98,7 +98,7 @@ struct invoke_policy_impl : invoke_policy_base<FilteredPattern> {
>::type >::type
mimpl; mimpl;
util::fixed_vector<size_t,util::tl_size<FilteredPattern>::value> mv; util::limited_vector<size_t,util::tl_size<FilteredPattern>::value> mv;
if (type_token == typeid(FilteredPattern) || mimpl::_(tup, mv)) { if (type_token == typeid(FilteredPattern) || mimpl::_(tup, mv)) {
for (size_t i = 0; i < util::tl_size<FilteredPattern>::value; ++i) { for (size_t i = 0; i < util::tl_size<FilteredPattern>::value; ++i) {
result[i] = const_cast<void*>(tup.at(mv[i])); result[i] = const_cast<void*>(tup.at(mv[i]));
...@@ -781,7 +781,7 @@ class match_expr { ...@@ -781,7 +781,7 @@ class match_expr {
typedef std::pair<const std::type_info*,std::uint64_t> cache_element; typedef std::pair<const std::type_info*,std::uint64_t> cache_element;
util::fixed_vector<cache_element,cache_size> m_cache; util::limited_vector<cache_element,cache_size> m_cache;
// ring buffer like access to m_cache // ring buffer like access to m_cache
size_t m_cache_begin; size_t m_cache_begin;
......
...@@ -33,7 +33,10 @@ ...@@ -33,7 +33,10 @@
#include "cppa/option.hpp" #include "cppa/option.hpp"
#include "cppa/cow_tuple.hpp" #include "cppa/cow_tuple.hpp"
#include "cppa/util/call.hpp" #include "cppa/util/call.hpp"
#include "cppa/util/limited_vector.hpp"
#include "cppa/opencl/global.hpp" #include "cppa/opencl/global.hpp"
#include "cppa/opencl/command_dispatcher.hpp" #include "cppa/opencl/command_dispatcher.hpp"
...@@ -70,11 +73,11 @@ struct cl_spawn_helper<std::function<option<cow_tuple<Ts...>> (any_tuple)>, ...@@ -70,11 +73,11 @@ struct cl_spawn_helper<std::function<option<cow_tuple<Ts...>> (any_tuple)>,
template<typename Signature, typename... Ts> template<typename Signature, typename... Ts>
inline actor_ptr spawn_cl(const opencl::program& prog, inline actor_ptr spawn_cl(const opencl::program& prog,
const char* fname, const char* fname,
std::vector<size_t> dims, const opencl::dim_vec& dims,
std::vector<size_t> offset = {}, const opencl::dim_vec& offset = {},
std::vector<size_t> local_dims = {}) { const opencl::dim_vec& local_dims = {}) {
detail::cl_spawn_helper<Signature> f; detail::cl_spawn_helper<Signature> f;
return util::call_mv(f, prog, fname, dims, offset, local_dims); return f(prog, fname, dims, offset, local_dims);
} }
/** /**
...@@ -87,12 +90,12 @@ inline actor_ptr spawn_cl(const opencl::program& prog, ...@@ -87,12 +90,12 @@ inline actor_ptr spawn_cl(const opencl::program& prog,
template<typename Signature, typename... Ts> template<typename Signature, typename... Ts>
inline actor_ptr spawn_cl(const char* source, inline actor_ptr spawn_cl(const char* source,
const char* fname, const char* fname,
std::vector<size_t> dims, const opencl::dim_vec& dims,
std::vector<size_t> offset = {}, const opencl::dim_vec& offset = {},
std::vector<size_t> local_dims = {}) { const opencl::dim_vec& local_dims = {}) {
auto prog = opencl::program::create(source); auto prog = opencl::program::create(source);
detail::cl_spawn_helper<Signature> f; detail::cl_spawn_helper<Signature> f;
return util::call_mv(f, prog, fname, dims, offset, local_dims); return f(prog, fname, dims, offset, local_dims);
} }
/** /**
...@@ -109,13 +112,13 @@ inline actor_ptr spawn_cl(const opencl::program& prog, ...@@ -109,13 +112,13 @@ inline actor_ptr spawn_cl(const opencl::program& prog,
const char* fname, const char* fname,
MapArgs map_args, MapArgs map_args,
MapResult map_result, MapResult map_result,
std::vector<size_t> dims, const opencl::dim_vec& dims,
std::vector<size_t> offset = {}, const opencl::dim_vec& offset = {},
std::vector<size_t> local_dims = {}) { const opencl::dim_vec& local_dims = {}) {
typedef typename util::get_callable_trait<MapArgs>::fun_type f0; typedef typename util::get_callable_trait<MapArgs>::fun_type f0;
typedef typename util::get_callable_trait<MapResult>::fun_type f1; typedef typename util::get_callable_trait<MapResult>::fun_type f1;
detail::cl_spawn_helper<f0,f1> f; detail::cl_spawn_helper<f0,f1> f;
return util::call_mv(f, prog, fname, dims, offset, local_dims, return f(prog, fname, dims, offset, local_dims,
f0{map_args}, f1{map_result}); f0{map_args}, f1{map_result});
} }
...@@ -133,13 +136,13 @@ inline actor_ptr spawn_cl(const char* source, ...@@ -133,13 +136,13 @@ inline actor_ptr spawn_cl(const char* source,
const char* fun_name, const char* fun_name,
MapArgs map_args, MapArgs map_args,
MapResult map_result, MapResult map_result,
std::vector<size_t> dims, const opencl::dim_vec& dims,
std::vector<size_t> offset = {}, const opencl::dim_vec& offset = {},
std::vector<size_t> local_dims = {}) { const opencl::dim_vec& local_dims = {}) {
using std::move; using std::move;
return spawn_cl(opencl::program::create(source), fun_name, return spawn_cl(opencl::program::create(source), fun_name,
move(map_args), move(map_result), move(map_args), move(map_result),
move(dims), move(offset), move(local_dims)); dims, offset, local_dims);
} }
} // namespace cppa } // namespace cppa
......
...@@ -45,7 +45,7 @@ ...@@ -45,7 +45,7 @@
#include "cppa/util/int_list.hpp" #include "cppa/util/int_list.hpp"
#include "cppa/util/type_list.hpp" #include "cppa/util/type_list.hpp"
#include "cppa/util/scope_guard.hpp" #include "cppa/util/limited_vector.hpp"
#include "cppa/opencl/global.hpp" #include "cppa/opencl/global.hpp"
#include "cppa/opencl/command.hpp" #include "cppa/opencl/command.hpp"
...@@ -75,20 +75,27 @@ class actor_facade<Ret(Args...)> : public actor { ...@@ -75,20 +75,27 @@ class actor_facade<Ret(Args...)> : public actor {
static actor_facade* create(command_dispatcher* dispatcher, static actor_facade* create(command_dispatcher* dispatcher,
const program& prog, const program& prog,
const char* kernel_name, const char* kernel_name,
std::vector<size_t> global_dimensions, const dim_vec& global_dims,
std::vector<size_t> global_offsets, const dim_vec& offsets,
std::vector<size_t> local_dimensions, const dim_vec& local_dims,
arg_mapping map_args, arg_mapping map_args,
result_mapping map_result) { result_mapping map_result) {
if(local_dimensions .size() > 3 || if (global_dims.empty()) {
global_dimensions.size() > 3 || auto str = "OpenCL kernel needs at least 1 global dimension.";
global_dimensions.empty()) { CPPA_LOGM_ERROR(detail::demangle(typeid(actor_facade)), str);
throw std::runtime_error(str);
}
auto check_vec = [&](const dim_vec& vec, const char* name) {
if (!vec.empty() && vec.size() != global_dims.size()) {
std::ostringstream oss; std::ostringstream oss;
oss << "OpenCL kernel allows a maximum of 3 dimensions" oss << name << " vector is not empty, but "
" and needs at least 1 global dimension."; << "its size differs from global dimensions vector's size";
CPPA_LOGM_ERROR(detail::demangle(typeid(actor_facade)), oss.str()); CPPA_LOGM_ERROR(detail::demangle<actor_facade>(), oss.str());
throw std::runtime_error(oss.str()); throw std::runtime_error(oss.str());
} }
};
check_vec(offsets, "offsets");
check_vec(local_dims, "local dimensions");
cl_int err{0}; cl_int err{0};
kernel_ptr kernel; kernel_ptr kernel;
kernel.adopt(clCreateKernel(prog.m_program.get(), kernel.adopt(clCreateKernel(prog.m_program.get(),
...@@ -105,9 +112,9 @@ class actor_facade<Ret(Args...)> : public actor { ...@@ -105,9 +112,9 @@ class actor_facade<Ret(Args...)> : public actor {
return new actor_facade<Ret (Args...)>{dispatcher, return new actor_facade<Ret (Args...)>{dispatcher,
kernel, kernel,
prog, prog,
std::move(global_dimensions), std::move(global_dims),
std::move(global_offsets), std::move(offsets),
std::move(local_dimensions), std::move(local_dims),
std::move(map_args), std::move(map_args),
std::move(map_result)}; std::move(map_result)};
} }
...@@ -130,18 +137,18 @@ class actor_facade<Ret(Args...)> : public actor { ...@@ -130,18 +137,18 @@ class actor_facade<Ret(Args...)> : public actor {
actor_facade(command_dispatcher* dispatcher, actor_facade(command_dispatcher* dispatcher,
kernel_ptr kernel, kernel_ptr kernel,
const program& prog, const program& prog,
std::vector<size_t> global_dimensions, const dim_vec& global_dimensions,
std::vector<size_t> global_offsets, const dim_vec& global_offsets,
std::vector<size_t> local_dimensions, const dim_vec& local_dimensions,
arg_mapping map_args, arg_mapping map_args,
result_mapping map_result) result_mapping map_result)
: m_kernel(kernel) : m_kernel(kernel)
, m_program(prog.m_program) , m_program(prog.m_program)
, m_context(prog.m_context) , m_context(prog.m_context)
, m_dispatcher(dispatcher) , m_dispatcher(dispatcher)
, m_global_dimensions(std::move(global_dimensions)) , m_global_dimensions(global_dimensions)
, m_global_offsets(std::move(global_offsets)) , m_global_offsets(global_offsets)
, m_local_dimensions(std::move(local_dimensions)) , m_local_dimensions(local_dimensions)
, m_map_args(std::move(map_args)) , m_map_args(std::move(map_args))
, m_map_result(std::move(map_result)) , m_map_result(std::move(map_result))
{ {
...@@ -185,9 +192,9 @@ class actor_facade<Ret(Args...)> : public actor { ...@@ -185,9 +192,9 @@ class actor_facade<Ret(Args...)> : public actor {
program_ptr m_program; program_ptr m_program;
context_ptr m_context; context_ptr m_context;
command_dispatcher* m_dispatcher; command_dispatcher* m_dispatcher;
std::vector<size_t> m_global_dimensions; dim_vec m_global_dimensions;
std::vector<size_t> m_global_offsets; dim_vec m_global_offsets;
std::vector<size_t> m_local_dimensions; dim_vec m_local_dimensions;
arg_mapping m_map_args; arg_mapping m_map_args;
result_mapping m_map_result; result_mapping m_map_result;
......
...@@ -70,9 +70,9 @@ class command_impl : public command { ...@@ -70,9 +70,9 @@ class command_impl : public command {
command_impl(response_handle handle, command_impl(response_handle handle,
kernel_ptr kernel, kernel_ptr kernel,
std::vector<mem_ptr> arguments, std::vector<mem_ptr> arguments,
const std::vector<size_t>& global_dimensions, const dim_vec& global_dimensions,
const std::vector<size_t>& global_offsets, const dim_vec& global_offsets,
const std::vector<size_t>& local_dimensions, const dim_vec& local_dimensions,
const std::function<any_tuple(T&)>& map_result) const std::function<any_tuple(T&)>& map_result)
: m_number_of_values(1) : m_number_of_values(1)
, m_handle(handle) , m_handle(handle)
...@@ -84,9 +84,7 @@ class command_impl : public command { ...@@ -84,9 +84,7 @@ class command_impl : public command {
, m_map_result(map_result) , m_map_result(map_result)
{ {
m_kernel_event.adopt(cl_event()); m_kernel_event.adopt(cl_event());
std::for_each(m_global_dimensions.begin(), for (auto s : m_global_dimensions) m_number_of_values *= s;
m_global_dimensions.end(),
[&](const size_t& s) { m_number_of_values *= s; });
} }
void enqueue (command_queue_ptr queue) { void enqueue (command_queue_ptr queue) {
...@@ -97,20 +95,23 @@ class command_impl : public command { ...@@ -97,20 +95,23 @@ class command_impl : public command {
auto ptr = m_kernel_event.get(); auto ptr = m_kernel_event.get();
auto data_or_nullptr = [](const dim_vec& vec) {
return vec.empty() ? nullptr : vec.data();
};
/* enqueue kernel */ /* enqueue kernel */
err = clEnqueueNDRangeKernel(m_queue.get(), err = clEnqueueNDRangeKernel(m_queue.get(),
m_kernel.get(), m_kernel.get(),
m_global_dimensions.size(), m_global_dimensions.size(),
m_global_offsets.data(), data_or_nullptr(m_global_offsets),
m_global_dimensions.data(), data_or_nullptr(m_global_dimensions),
m_local_dimensions.data(), data_or_nullptr(m_local_dimensions),
0, 0,
nullptr, nullptr,
&ptr); &ptr);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clEnqueueNDRangeKernel: '" throw std::runtime_error("clEnqueueNDRangeKernel: "
+ get_opencl_error(err) + get_opencl_error(err));
+ "'.");
} }
err = clSetEventCallback(ptr, err = clSetEventCallback(ptr,
CL_COMPLETE, CL_COMPLETE,
...@@ -124,9 +125,8 @@ class command_impl : public command { ...@@ -124,9 +125,8 @@ class command_impl : public command {
}, },
this); this);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clSetEventCallback: '" throw std::runtime_error("clSetEventCallback: "
+ get_opencl_error(err) + get_opencl_error(err));
+ "'.");
} }
} }
...@@ -138,9 +138,9 @@ class command_impl : public command { ...@@ -138,9 +138,9 @@ class command_impl : public command {
event_ptr m_kernel_event; event_ptr m_kernel_event;
command_queue_ptr m_queue; command_queue_ptr m_queue;
std::vector<mem_ptr> m_arguments; std::vector<mem_ptr> m_arguments;
std::vector<size_t> m_global_dimensions; dim_vec m_global_dimensions;
std::vector<size_t> m_global_offsets; dim_vec m_global_offsets;
std::vector<size_t> m_local_dimensions; dim_vec m_local_dimensions;
std::function<any_tuple (T&)> m_map_result; std::function<any_tuple (T&)> m_map_result;
void handle_results () { void handle_results () {
...@@ -166,7 +166,6 @@ class command_impl : public command { ...@@ -166,7 +166,6 @@ class command_impl : public command {
} }
auto mapped_result = m_map_result(result); auto mapped_result = m_map_result(result);
reply_tuple_to(m_handle, mapped_result); reply_tuple_to(m_handle, mapped_result);
//reply_to(m_handle, results);
} }
}; };
......
...@@ -38,9 +38,9 @@ ...@@ -38,9 +38,9 @@
#include <functional> #include <functional>
#include "cppa/option.hpp" #include "cppa/option.hpp"
#include "cppa/logging.hpp"
#include "cppa/channel.hpp" #include "cppa/channel.hpp"
#include "cppa/cow_tuple.hpp" #include "cppa/cow_tuple.hpp"
#include "cppa/logging.hpp"
#include "cppa/opencl/global.hpp" #include "cppa/opencl/global.hpp"
#include "cppa/opencl/command.hpp" #include "cppa/opencl/command.hpp"
...@@ -48,6 +48,8 @@ ...@@ -48,6 +48,8 @@
#include "cppa/opencl/smart_ptr.hpp" #include "cppa/opencl/smart_ptr.hpp"
#include "cppa/opencl/actor_facade.hpp" #include "cppa/opencl/actor_facade.hpp"
#include "cppa/util/limited_vector.hpp"
#include "cppa/detail/singleton_mixin.hpp" #include "cppa/detail/singleton_mixin.hpp"
#include "cppa/detail/singleton_manager.hpp" #include "cppa/detail/singleton_manager.hpp"
...@@ -84,18 +86,18 @@ class command_dispatcher { ...@@ -84,18 +86,18 @@ class command_dispatcher {
template<typename Ret, typename... Args> template<typename Ret, typename... Args>
actor_ptr spawn(const program& prog, actor_ptr spawn(const program& prog,
const char* kernel_name, const char* kernel_name,
std::vector<size_t> global_dims, const dim_vec& global_dims,
std::vector<size_t> global_offs, const dim_vec& global_offs,
std::vector<size_t> local_dims, const dim_vec& local_dims,
std::function<option<cow_tuple<typename util::rm_ref<Args>::type...>>(any_tuple)> map_args, std::function<option<cow_tuple<typename util::rm_ref<Args>::type...>>(any_tuple)> map_args,
std::function<any_tuple(Ret&)> map_result) std::function<any_tuple(Ret&)> map_result)
{ {
return actor_facade<Ret (Args...)>::create(this, return actor_facade<Ret (Args...)>::create(this,
prog, prog,
kernel_name, kernel_name,
std::move(global_dims), global_dims,
std::move(global_offs), global_offs,
std::move(local_dims), local_dims,
std::move(map_args), std::move(map_args),
std::move(map_result)); std::move(map_result));
} }
...@@ -103,9 +105,9 @@ class command_dispatcher { ...@@ -103,9 +105,9 @@ class command_dispatcher {
template<typename Ret, typename... Args> template<typename Ret, typename... Args>
actor_ptr spawn(const program& prog, actor_ptr spawn(const program& prog,
const char* kernel_name, const char* kernel_name,
std::vector<size_t> global_dims, const dim_vec& global_dims,
std::vector<size_t> global_offs = {}, const dim_vec& global_offs = {},
std::vector<size_t> local_dims = {}) const dim_vec& local_dims = {})
{ {
std::function<option<cow_tuple<typename util::rm_ref<Args>::type...>>(any_tuple)> std::function<option<cow_tuple<typename util::rm_ref<Args>::type...>>(any_tuple)>
map_args = [] (any_tuple msg) { map_args = [] (any_tuple msg) {
...@@ -131,20 +133,20 @@ class command_dispatcher { ...@@ -131,20 +133,20 @@ class command_dispatcher {
device_ptr dev_id; device_ptr dev_id;
size_t max_itms_per_grp; size_t max_itms_per_grp;
cl_uint max_dim; cl_uint max_dim;
std::vector<size_t> max_itms_per_dim; dim_vec max_itms_per_dim;
device_info(unsigned id, device_info(unsigned id,
command_queue_ptr queue, command_queue_ptr queue,
device_ptr device_id, device_ptr device_id,
size_t max_itms_per_grp, size_t max_itms_per_grp,
cl_uint max_dim, cl_uint max_dim,
std::vector<size_t> max_itms_per_dim) const dim_vec& max_itms_per_dim)
: id(id) : id(id)
, cmd_queue(queue) , cmd_queue(queue)
, dev_id(device_id) , dev_id(device_id)
, max_itms_per_grp(max_itms_per_grp) , max_itms_per_grp(max_itms_per_grp)
, max_dim(max_dim) , max_dim(max_dim)
, max_itms_per_dim(std::move(max_itms_per_dim)) { } , max_itms_per_dim(max_itms_per_dim) { }
}; };
typedef intrusive::blocking_single_reader_queue<command,dereferencer> typedef intrusive::blocking_single_reader_queue<command,dereferencer>
......
...@@ -34,6 +34,8 @@ ...@@ -34,6 +34,8 @@
#include <string> #include <string>
#include "cppa/util/limited_vector.hpp"
#if defined __APPLE__ || defined(MACOSX) #if defined __APPLE__ || defined(MACOSX)
#include <OpenCL/opencl.h> #include <OpenCL/opencl.h>
#else #else
...@@ -42,6 +44,11 @@ ...@@ -42,6 +44,11 @@
namespace cppa { namespace opencl { namespace cppa { namespace opencl {
/**
* @brief A vector of up to three elements used for OpenCL dimensions.
*/
typedef util::limited_vector<size_t,3> dim_vec;
std::string get_opencl_error(cl_int err); std::string get_opencl_error(cl_int err);
cl_int clReleaseDeviceDummy (cl_device_id); cl_int clReleaseDeviceDummy (cl_device_id);
......
...@@ -47,7 +47,7 @@ namespace cppa { namespace util { ...@@ -47,7 +47,7 @@ namespace cppa { namespace util {
* does <b>not</b> call constructors or destructors properly. * does <b>not</b> call constructors or destructors properly.
*/ */
template<typename T, size_t MaxSize> template<typename T, size_t MaxSize>
class fixed_vector { class limited_vector {
//static_assert(std::is_arithmetic<T>::value || std::is_pointer<T>::value, //static_assert(std::is_arithmetic<T>::value || std::is_pointer<T>::value,
// "T must be an arithmetic or pointer type"); // "T must be an arithmetic or pointer type");
...@@ -72,13 +72,17 @@ class fixed_vector { ...@@ -72,13 +72,17 @@ class fixed_vector {
typedef std::reverse_iterator<iterator> reverse_iterator; typedef std::reverse_iterator<iterator> reverse_iterator;
typedef std::reverse_iterator<const_iterator> const_reverse_iterator; typedef std::reverse_iterator<const_iterator> const_reverse_iterator;
inline fixed_vector() : m_size(0) { } inline limited_vector() : m_size(0) { }
fixed_vector(const fixed_vector& other) : m_size(other.m_size) { limited_vector(size_t initial_size) : m_size(initial_size) {
std::fill_n(begin(), initial_size, 0);
}
limited_vector(const limited_vector& other) : m_size(other.m_size) {
std::copy(other.begin(), other.end(), begin()); std::copy(other.begin(), other.end(), begin());
} }
fixed_vector& operator=(const fixed_vector& other) { limited_vector& operator=(const limited_vector& other) {
resize(other.size()); resize(other.size());
std::copy(other.begin(), other.end(), begin()); std::copy(other.begin(), other.end(), begin());
return *this; return *this;
...@@ -89,7 +93,7 @@ class fixed_vector { ...@@ -89,7 +93,7 @@ class fixed_vector {
m_size = s; m_size = s;
} }
fixed_vector(std::initializer_list<T> init) : m_size(init.size()) { limited_vector(std::initializer_list<T> init) : m_size(init.size()) {
CPPA_REQUIRE(init.size() <= MaxSize); CPPA_REQUIRE(init.size() <= MaxSize);
std::copy(init.begin(), init.end(), begin()); std::copy(init.begin(), init.end(), begin());
} }
...@@ -244,7 +248,7 @@ class fixed_vector { ...@@ -244,7 +248,7 @@ class fixed_vector {
inline void insert(iterator pos, InputIterator first, InputIterator last) { inline void insert(iterator pos, InputIterator first, InputIterator last) {
auto num_elements = std::distance(first, last); auto num_elements = std::distance(first, last);
if ((size() + num_elements) > MaxSize) { if ((size() + num_elements) > MaxSize) {
throw std::length_error("fixed_vector::insert: too much elements"); throw std::length_error("limited_vector::insert: too much elements");
} }
if (pos == end()) { if (pos == end()) {
resize(size() + num_elements); resize(size() + num_elements);
......
...@@ -241,7 +241,7 @@ void command_dispatcher::initialize() { ...@@ -241,7 +241,7 @@ void command_dispatcher::initialize() {
CPPA_LOG_ERROR(oss.str()); CPPA_LOG_ERROR(oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
vector<size_t> max_work_items_per_dim(max_work_item_dimensions); dim_vec max_work_items_per_dim(max_work_item_dimensions);
err = clGetDeviceInfo(device.get(), err = clGetDeviceInfo(device.get(),
CL_DEVICE_MAX_WORK_ITEM_SIZES, CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(size_t)*max_work_item_dimensions, sizeof(size_t)*max_work_item_dimensions,
...@@ -261,7 +261,7 @@ void command_dispatcher::initialize() { ...@@ -261,7 +261,7 @@ void command_dispatcher::initialize() {
device, device,
max_work_group_size, max_work_group_size,
max_work_item_dimensions, max_work_item_dimensions,
move(max_work_items_per_dim)}; max_work_items_per_dim};
m_devices.push_back(move(dev_info)); m_devices.push_back(move(dev_info));
} }
} }
......
...@@ -4,22 +4,22 @@ ...@@ -4,22 +4,22 @@
#include <algorithm> #include <algorithm>
#include "test.hpp" #include "test.hpp"
#include "cppa/util/fixed_vector.hpp" #include "cppa/util/limited_vector.hpp"
using std::cout; using std::cout;
using std::endl; using std::endl;
using std::equal; using std::equal;
using cppa::util::fixed_vector; using cppa::util::limited_vector;
int main() { int main() {
CPPA_TEST(test_fixed_vector); CPPA_TEST(test_limited_vector);
int arr1[] {1, 2, 3, 4}; int arr1[] {1, 2, 3, 4};
fixed_vector<int,4> vec1 {1, 2, 3, 4}; limited_vector<int,4> vec1 {1, 2, 3, 4};
fixed_vector<int,5> vec2 {4, 3, 2, 1}; limited_vector<int,5> vec2 {4, 3, 2, 1};
fixed_vector<int,4> vec3; limited_vector<int,4> vec3;
for (int i = 1; i <= 4; ++i) vec3.push_back(i); for (int i = 1; i <= 4; ++i) vec3.push_back(i);
fixed_vector<int,4> vec4 {1, 2}; limited_vector<int,4> vec4 {1, 2};
fixed_vector<int,2> vec5 {3, 4}; limited_vector<int,2> vec5 {3, 4};
vec4.insert(vec4.end(), vec5.begin(), vec5.end()); vec4.insert(vec4.end(), vec5.begin(), vec5.end());
auto vec6 = vec4; auto vec6 = vec4;
CPPA_CHECK_EQUAL(vec1.size(), 4); CPPA_CHECK_EQUAL(vec1.size(), 4);
...@@ -39,13 +39,13 @@ int main() { ...@@ -39,13 +39,13 @@ int main() {
CPPA_CHECK(std::equal(vec4.begin(), vec4.end(), arr1)); CPPA_CHECK(std::equal(vec4.begin(), vec4.end(), arr1));
CPPA_CHECK(std::equal(vec6.begin(), vec6.end(), arr1)); CPPA_CHECK(std::equal(vec6.begin(), vec6.end(), arr1));
CPPA_CHECK(std::equal(vec6.begin(), vec6.end(), vec2.rbegin())); CPPA_CHECK(std::equal(vec6.begin(), vec6.end(), vec2.rbegin()));
fixed_vector<int,10> vec7 {5, 9}; limited_vector<int,10> vec7 {5, 9};
fixed_vector<int,10> vec8 {1, 2, 3, 4}; limited_vector<int,10> vec8 {1, 2, 3, 4};
fixed_vector<int,10> vec9 {6, 7, 8}; limited_vector<int,10> vec9 {6, 7, 8};
vec7.insert(vec7.begin() + 1, vec9.begin(), vec9.end()); vec7.insert(vec7.begin() + 1, vec9.begin(), vec9.end());
vec7.insert(vec7.begin(), vec8.begin(), vec8.end()); vec7.insert(vec7.begin(), vec8.begin(), vec8.end());
CPPA_CHECK_EQUAL(vec7.full(), false); CPPA_CHECK_EQUAL(vec7.full(), false);
fixed_vector<int,1> vec10 {10}; limited_vector<int,1> vec10 {10};
vec7.insert(vec7.end(), vec10.begin(), vec10.end()); vec7.insert(vec7.end(), vec10.begin(), vec10.end());
CPPA_CHECK_EQUAL(vec7.full(), true); CPPA_CHECK_EQUAL(vec7.full(), true);
CPPA_CHECK((std::is_sorted(vec7.begin(), vec7.end()))); CPPA_CHECK((std::is_sorted(vec7.begin(), vec7.end())));
......
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment