Commit a828a3c7 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 76c023e1
...@@ -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);
......
...@@ -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));
} }
} }
......
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