Commit 7d1ea179 authored by Joseph Noir's avatar Joseph Noir

Update syntax of spawn_cl

The new syntax describes the signature in form of a variadic number of
in, out and in_out arguments as the last arguments to spawn_cl. This
allows multiple output arguments as well as arguments used for input as
well as output. The type of the argument is passed as a template
parameter to the in, out and in_out classes.

out accepts a function to calculate the size of the result buffer
from the input arguments for a specific calculation. Per default, its
size is equal to the number of work items.

The configuration parameters for the actor (work-items, offsets, ...)
are now wrapped in a spawn_configuration.

Examples and tests have been updated accordingly. For now, the old
spawn_cl syntax is still valid, but marked as deprecated.
parent 626b3456
......@@ -38,43 +38,72 @@
#include "caf/opencl/global.hpp"
#include "caf/opencl/command.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/spawn_config.hpp"
namespace caf {
namespace opencl {
class opencl_metainfo;
template <class Signature>
class actor_facade;
template <class List>
struct function_sig_from_outputs;
template <class Ret, typename... Args>
class actor_facade<Ret(Args...)> : public abstract_actor {
template <class... Ts>
struct function_sig_from_outputs<detail::type_list<Ts...>> {
using type = std::function<message (Ts...)>;
};
template <class T, class List>
struct command_sig_from_outputs;
friend class command<actor_facade, Ret>;
template <class T, class... Ts>
struct command_sig_from_outputs<T, detail::type_list<Ts...>> {
using type = command<T, Ts...>;
};
template <class... Ts>
class actor_facade : public abstract_actor {
public:
using arg_types = detail::type_list<typename std::decay<Args>::type...>;
using arg_mapping = std::function<optional<message> (message&)>;
using result_mapping = std::function<message(Ret&)>;
using arg_types = detail::type_list<Ts...>;
using unpacked_types = typename detail::tl_map<arg_types, extract_type>::type;
using input_wrapped_types =
typename detail::tl_filter<arg_types, is_input_arg>::type;
using input_types =
typename detail::tl_map<input_wrapped_types, extract_type>::type;
using input_mapping = std::function<optional<message> (message&)>;
using output_wrapped_types =
typename detail::tl_filter<arg_types, is_output_arg>::type;
using output_types =
typename detail::tl_map<output_wrapped_types, extract_type>::type;
using output_mapping = typename function_sig_from_outputs<output_types>::type;
typename detail::il_indices<arg_types>::type indices;
using evnt_vec = std::vector<cl_event>;
using args_vec = std::vector<mem_ptr>;
using command_type = command<actor_facade, Ret>;
static intrusive_ptr<actor_facade>
create(const program& prog, const char* kernel_name,
const dim_vec& global_dims, const dim_vec& offsets,
const dim_vec& local_dims, size_t result_size,
arg_mapping map_args = arg_mapping{},
result_mapping map_result = result_mapping{}) {
if (global_dims.empty()) {
using size_vec = std::vector<size_t>;
using command_type =
typename command_sig_from_outputs<actor_facade, output_types>::type;
static intrusive_ptr<actor_facade> create(const program& prog,
const char* kernel_name,
const spawn_config& config,
input_mapping map_args,
output_mapping map_result,
Ts&&... xs) {
if (config.dimensions().empty()) {
auto str = "OpenCL kernel needs at least 1 global dimension.";
CAF_LOGF_ERROR(str);
throw std::runtime_error(str);
}
auto check_vec = [&](const dim_vec& vec, const char* name) {
if (! vec.empty() && vec.size() != global_dims.size()) {
if (! vec.empty() && vec.size() != config.dimensions().size()) {
std::ostringstream oss;
oss << name << " vector is not empty, but "
<< "its size differs from global dimensions vector's size";
......@@ -82,19 +111,16 @@ public:
throw std::runtime_error(oss.str());
}
};
check_vec(offsets, "offsets");
check_vec(local_dims, "local dimensions");
kernel_ptr kernel;
kernel.reset(v2get(CAF_CLF(clCreateKernel), prog.program_.get(),
kernel_name),
false);
if (result_size == 0) {
result_size = std::accumulate(global_dims.begin(), global_dims.end(),
size_t{1}, std::multiplies<size_t>{});
check_vec(config.offsets(), "offsets");
check_vec(config.local_dimensions(), "local dimensions");
auto itr = prog.available_kernels_.find(kernel_name);
if (itr == prog.available_kernels_.end()) {
return nullptr;
} else {
return new actor_facade(prog, itr->second, config,
std::move(map_args), std::move(map_result),
std::forward_as_tuple(xs...));
}
return new actor_facade(prog, kernel, global_dims, offsets, local_dims,
result_size, std::move(map_args),
std::move(map_result));
}
void enqueue(const actor_addr &sender, message_id mid, message content,
......@@ -107,94 +133,137 @@ public:
}
content = std::move(*mapped);
}
typename detail::il_indices<arg_types>::type indices;
if (! content.match_elements(arg_types{})) {
if (! content.match_elements(input_types{})) {
return;
}
response_promise handle{this->address(), sender, mid.response_id()};
evnt_vec events;
args_vec arguments;
add_arguments_to_kernel<Ret>(events, arguments, result_size_,
content, indices);
args_vec input_buffers;
args_vec output_buffers;
size_vec result_sizes;
add_kernel_arguments(events, input_buffers, output_buffers,
result_sizes, content, indices);
auto cmd = make_counted<command_type>(handle, this,
std::move(events),
std::move(arguments),
result_size_,
std::move(input_buffers),
std::move(output_buffers),
std::move(result_sizes),
std::move(content));
cmd->enqueue();
}
private:
actor_facade(const program& prog, kernel_ptr kernel,
const dim_vec& global_dimensions, const dim_vec& global_offsets,
const dim_vec& local_dimensions, size_t result_size,
arg_mapping map_args, result_mapping map_result)
const spawn_config& config,
input_mapping map_args, output_mapping map_result,
std::tuple<Ts...> xs)
: kernel_(kernel),
program_(prog.program_),
context_(prog.context_),
queue_(prog.queue_),
global_dimensions_(global_dimensions),
global_offsets_(global_offsets),
local_dimensions_(local_dimensions),
result_size_(result_size),
config_(config),
map_args_(std::move(map_args)),
map_result_(std::move(map_result)) {
map_results_(std::move(map_result)),
argument_types_(xs) {
CAF_LOG_TRACE("id: " << this->id());
default_output_size_ = std::accumulate(config_.dimensions().begin(),
config_.dimensions().end(),
size_t{1},
std::multiplies<size_t>{});
}
void add_arguments_to_kernel_rec(evnt_vec&, args_vec& arguments, message&,
detail::int_list<>) {
// rotate left (output buffer to the end)
std::rotate(arguments.begin(), arguments.begin() + 1, arguments.end());
for (cl_uint i = 0; i < arguments.size(); ++i) {
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), i,
sizeof(cl_mem), static_cast<void*>(&arguments[i]));
}
clFlush(queue_.get());
void add_kernel_arguments(evnt_vec&, args_vec&, args_vec&, size_vec&,
message&, detail::int_list<>) {
// nop
}
/// the separation into input and output is required, because we need to
/// access the output arguments later on, but only keep the ptrs to the input
/// to prevent them from being deleted before our operation finished
template <long I, long... Is>
void add_arguments_to_kernel_rec(evnt_vec& events, args_vec& arguments,
void add_kernel_arguments(evnt_vec& events, args_vec& input_buffers,
args_vec& output_buffers, size_vec& sizes,
message& msg, detail::int_list<I, Is...>) {
using value_type = typename detail::tl_at<arg_types, I>::type;
auto& arg = msg.get_as<value_type>(I);
size_t buffer_size = sizeof(value_type) * arg.size();
create_buffer<I>(std::get<I>(argument_types_), events, sizes,
input_buffers, output_buffers, msg);
add_kernel_arguments(events, input_buffers, output_buffers, sizes, msg,
detail::int_list<Is...>{});
}
template <long I, class T>
void create_buffer(const in<T>&, evnt_vec& events, size_vec&,
args_vec& input_buffers, args_vec&, message& msg) {
using container_type = typename detail::tl_at<unpacked_types, I>::type;
using value_type = typename container_type::value_type;
auto& value = msg.get_as<container_type>(I);
auto size = value.size();
size_t buffer_size = sizeof(value_type) * size;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
cl_mem_flags{CL_MEM_READ_ONLY}, buffer_size, nullptr);
cl_event event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
cl_mem_flags{CL_MEM_READ_WRITE}, buffer_size, nullptr);
auto event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
queue_.get(), buffer, cl_bool{CL_FALSE},
cl_uint{0}, buffer_size, arg.data());
cl_uint{0}, buffer_size, value.data());
events.push_back(std::move(event));
mem_ptr tmp;
tmp.reset(buffer, false);
arguments.push_back(tmp);
add_arguments_to_kernel_rec(events, arguments, msg,
detail::int_list<Is...>{});
input_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), I,
sizeof(cl_mem), static_cast<void*>(&input_buffers.back()));
}
template <class R, class Token>
void add_arguments_to_kernel(evnt_vec& events, args_vec& arguments,
size_t ret_size, message& msg, Token tk) {
arguments.clear();
auto buf = v2get(CAF_CLF(clCreateBuffer), context_.get(),
cl_mem_flags{CL_MEM_WRITE_ONLY},
sizeof(typename R::value_type) * ret_size, nullptr);
template <long I, class T>
void create_buffer(const in_out<T>&, evnt_vec& events, size_vec& sizes,
args_vec&, args_vec& output_buffers, message& msg) {
using container_type = typename detail::tl_at<unpacked_types, I>::type;
using value_type = typename container_type::value_type;
auto& value = msg.get_as<container_type>(I);
auto size = value.size();
size_t buffer_size = sizeof(value_type) * size;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
cl_mem_flags{CL_MEM_READ_WRITE}, buffer_size, nullptr);
auto event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
queue_.get(), buffer, cl_bool{CL_FALSE},
cl_uint{0}, buffer_size, value.data());
events.push_back(std::move(event));
mem_ptr tmp;
tmp.reset(buf, false);
arguments.push_back(tmp);
add_arguments_to_kernel_rec(events, arguments, msg, tk);
tmp.reset(buffer, false);
output_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), I,
sizeof(cl_mem), static_cast<void*>(&output_buffers.back()));
sizes.push_back(size);
}
template <long I, class T>
void create_buffer(const out<T>& wrapper, evnt_vec&, size_vec& sizes,
args_vec&, args_vec& output_buffers, message& msg) {
using container_type = typename detail::tl_at<unpacked_types, I>::type;
using value_type = typename container_type::value_type;
auto size = get_size_for_argument(wrapper, msg, default_output_size_);
auto buffer_size = sizeof(value_type) * size;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
cl_mem_flags{CL_MEM_READ_WRITE}, buffer_size, nullptr);
mem_ptr tmp;
tmp.reset(buffer, false);
output_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), I,
sizeof(cl_mem), static_cast<void*>(&output_buffers.back()));
sizes.push_back(size);
}
template <class Fun>
size_t get_size_for_argument(Fun& f, message& m, size_t default_size) {
auto size = f(m);
return size && (*size > 0) ? *size : default_size;
}
kernel_ptr kernel_;
program_ptr program_;
context_ptr context_;
command_queue_ptr queue_;
dim_vec global_dimensions_;
dim_vec global_offsets_;
dim_vec local_dimensions_;
size_t result_size_;
arg_mapping map_args_;
result_mapping map_result_;
spawn_config config_;
input_mapping map_args_;
output_mapping map_results_;
std::tuple<Ts...> argument_types_;
size_t default_output_size_;
};
} // namespace opencl
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* Raphael Hiesgen <raphael.hiesgen (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_ARGUMENTS
#define CAF_OPENCL_ARGUMENTS
#include <functional>
#include <type_traits>
#include "caf/message.hpp"
#include "caf/optional.hpp"
namespace caf {
namespace opencl {
/// Use as a default way to calculate output size. 0 will be set to the number
/// of work items at runtime.
struct dummy_size_calculator {
template <class... Ts>
size_t operator()(Ts&&...) const {
return 0;
}
};
/// Mark an a spawn_cl template argument as input only
template <class Arg>
struct in {
using arg_type = typename std::decay<Arg>::type;
};
/// Mark an a spawn_cl template argument as input and output
template <class Arg>
struct in_out {
using arg_type = typename std::decay<Arg>::type;
};
template <class Arg>
struct out {
out() { }
template <class F>
out(F fun) {
fun_ = [fun](message& msg) -> optional<size_t> {
auto res = msg.apply(fun);
size_t result;
if (res) {
res->apply([&](size_t x) { result = x; });
return result;
}
return none;
};
}
optional<size_t> operator()(message& msg) const {
return fun_ ? fun_(msg) : 0;
}
std::function<optional<size_t> (message&)> fun_;
};
///Cconverts C arrays, i.e., pointers, to vectors.
template <class T>
struct carr_to_vec {
using type = T;
};
template <class T>
struct carr_to_vec<T*> {
using type = std::vector<T>;
};
/// Filter types for any argument type.
template <class T>
struct is_opencl_arg : std::false_type {};
template <class T>
struct is_opencl_arg<in<T>> : std::true_type {};
template <class T>
struct is_opencl_arg<in_out<T>> : std::true_type {};
template <class T>
struct is_opencl_arg<out<T>> : std::true_type {};
/// Filter type lists for input arguments, in and in_out.
template <class T>
struct is_input_arg : std::false_type {};
template <class T>
struct is_input_arg<in<T>> : std::true_type {};
template <class T>
struct is_input_arg<in_out<T>> : std::true_type {};
/// Filter type lists for output arguments, in_out and out.
template <class T>
struct is_output_arg : std::false_type {};
template <class T>
struct is_output_arg<out<T>> : std::true_type {};
template <class T>
struct is_output_arg<in_out<T>> : std::true_type {};
/// Filter for arguments that require size, in this case only out.
template <class T>
struct requires_size_arg : std::false_type {};
template <class T>
struct requires_size_arg<out<T>> : std::true_type {};
/// extract types
template <class T>
struct extract_type { };
template <class T>
struct extract_type<in<T>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type;
};
template <class T>
struct extract_type<in_out<T>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type;
};
template <class T>
struct extract_type<out<T>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type;
};
/// Create the return message from tuple arumgent
struct message_from_results {
template <class T, class... Ts>
message operator()(T& x, Ts&... xs) {
return make_message(std::move(x), std::move(xs)...);
}
template <class... Ts>
message operator()(std::tuple<Ts...>& values) {
return apply_args(*this, detail::get_indices(values), values);
}
};
/// Helpers for conversion in deprecated spawn functions
template <class T>
struct to_input_arg {
using type = in<T>;
};
template <class T>
struct to_output_arg {
using type = out<T>;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_ARGUMENTS
......@@ -20,41 +20,49 @@
#ifndef CAF_OPENCL_COMMAND_HPP
#define CAF_OPENCL_COMMAND_HPP
#include <tuple>
#include <vector>
#include <numeric>
#include <algorithm>
#include <functional>
#include "caf/detail/logging.hpp"
#include "caf/opencl/global.hpp"
#include "caf/abstract_actor.hpp"
#include "caf/response_promise.hpp"
#include "caf/detail/logging.hpp"
#include "caf/detail/scope_guard.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/detail/scope_guard.hpp"
namespace caf {
namespace opencl {
template <typename T, typename R>
template <class FacadeType, class... Ts>
class command : public ref_counted {
public:
command(response_promise handle, intrusive_ptr<T> actor_facade,
std::vector<cl_event> events, std::vector<mem_ptr> arguments,
size_t result_size, message msg)
: result_size_(result_size),
command(response_promise handle, intrusive_ptr<FacadeType> actor_facade,
std::vector<cl_event> events, std::vector<mem_ptr> input_buffers,
std::vector<mem_ptr> output_buffers, std::vector<size_t> result_sizes,
message msg)
: result_sizes_(result_sizes),
handle_(handle),
actor_facade_(actor_facade),
queue_(actor_facade->queue_),
events_(std::move(events)),
arguments_(std::move(arguments)),
result_(result_size),
mem_in_events_(std::move(events)),
input_buffers_(std::move(input_buffers)),
output_buffers_(std::move(output_buffers)),
msg_(msg) {
// nop
}
~command() {
for (auto& e : events_) {
for (auto& e : mem_in_events_) {
v1callcl(CAF_CLF(clReleaseEvent),e);
}
for (auto& e : mem_out_events_) {
v1callcl(CAF_CLF(clReleaseEvent),e);
}
}
......@@ -71,28 +79,35 @@ public:
// OpenCL expects cl_uint (unsigned int), hence the cast
cl_int err = clEnqueueNDRangeKernel(
queue_.get(), actor_facade_->kernel_.get(),
static_cast<cl_uint>(actor_facade_->global_dimensions_.size()),
data_or_nullptr(actor_facade_->global_offsets_),
data_or_nullptr(actor_facade_->global_dimensions_),
data_or_nullptr(actor_facade_->local_dimensions_),
static_cast<cl_uint>(events_.size()),
(events_.empty() ? nullptr : events_.data()), &event_k);
static_cast<cl_uint>(actor_facade_->config_.dimensions().size()),
data_or_nullptr(actor_facade_->config_.offsets()),
data_or_nullptr(actor_facade_->config_.dimensions()),
data_or_nullptr(actor_facade_->config_.local_dimensions()),
static_cast<cl_uint>(mem_in_events_.size()),
(mem_in_events_.empty() ? nullptr : mem_in_events_.data()), &event_k
);
if (err != CL_SUCCESS) {
CAF_LOGMF(CAF_ERROR, "clEnqueueNDRangeKernel: " << get_opencl_error(err));
this->deref(); // or can anything actually happen?
clReleaseEvent(event_k);
this->deref();
return;
} else {
cl_event event_r;
err =
clEnqueueReadBuffer(queue_.get(), arguments_.back().get(), CL_FALSE,
0, sizeof(typename R::value_type) * result_size_,
result_.data(), 1, &event_k, &event_r);
enqueue_read_buffers(event_k, detail::get_indices(result_buffers_));
cl_event marker;
#if defined(__APPLE__)
err = clEnqueueMarkerWithWaitList(queue_.get(), mem_out_events_.size(),
mem_out_events_.data(), &marker);
#else
err = clEnqueueMarker(queue_.get(), &marker);
#endif
if (err != CL_SUCCESS) {
this->deref(); // failed to enqueue command
throw std::runtime_error("clEnqueueReadBuffer: " +
get_opencl_error(err));
CAF_LOGMF(CAF_ERROR, "clSetEventCallback: " << get_opencl_error(err));
clReleaseEvent(marker);
clReleaseEvent(event_k);
this->deref(); // callback is not set
return;
}
err = clSetEventCallback(event_r, CL_COMPLETE,
err = clSetEventCallback(marker, CL_COMPLETE,
[](cl_event, cl_int, void* data) {
auto cmd = reinterpret_cast<command*>(data);
cmd->handle_results();
......@@ -101,32 +116,64 @@ public:
this);
if (err != CL_SUCCESS) {
CAF_LOGMF(CAF_ERROR, "clSetEventCallback: " << get_opencl_error(err));
clReleaseEvent(marker);
clReleaseEvent(event_k);
this->deref(); // callback is not set
return;
}
err = clFlush(queue_.get());
if (err != CL_SUCCESS) {
CAF_LOGMF(CAF_ERROR, "clFlush: " << get_opencl_error(err));
}
events_.push_back(std::move(event_k));
events_.push_back(std::move(event_r));
mem_out_events_.push_back(std::move(event_k));
mem_out_events_.push_back(std::move(marker));
}
}
private:
size_t result_size_;
std::vector<size_t> result_sizes_;
response_promise handle_;
intrusive_ptr<T> actor_facade_;
intrusive_ptr<FacadeType> actor_facade_;
command_queue_ptr queue_;
std::vector<cl_event> events_;
std::vector<mem_ptr> arguments_;
R result_;
std::vector<cl_event> mem_in_events_;
std::vector<cl_event> mem_out_events_;
std::vector<mem_ptr> input_buffers_;
std::vector<mem_ptr> output_buffers_;
std::tuple<Ts...> result_buffers_;
message msg_; // required to keep the argument buffers alive (async copy)
void enqueue_read_buffers(cl_event&, detail::int_list<>) {
// nop, end of recursion
}
template <long I, long... Is>
void enqueue_read_buffers(cl_event& kernel_done, detail::int_list<I, Is...>) {
using container_type =
typename std::tuple_element<I, std::tuple<Ts...>>::type;
using value_type = typename container_type::value_type;
cl_event event;
auto size = result_sizes_[I];
auto buffer_size = sizeof(value_type) * result_sizes_[I];
std::get<I>(result_buffers_).resize(size);
auto err = clEnqueueReadBuffer(queue_.get(), output_buffers_[I].get(),
CL_FALSE, 0, buffer_size,
std::get<I>(result_buffers_).data(),
1, &kernel_done, &event);
if (err != CL_SUCCESS) {
this->deref(); // failed to enqueue command
throw std::runtime_error("clEnqueueReadBuffer: " +
get_opencl_error(err));
}
mem_out_events_.push_back(std::move(event));
enqueue_read_buffers(kernel_done, detail::int_list<Is...>{});
}
void handle_results() {
auto& map_fun = actor_facade_->map_result_;
auto msg = map_fun ? map_fun(result_) : make_message(std::move(result_));
auto& map_fun = actor_facade_->map_results_;
auto msg = map_fun ? apply_args(map_fun,
detail::get_indices(result_buffers_),
result_buffers_)
: message_from_results{}(result_buffers_);
handle_.deliver(std::move(msg));
}
};
......
......@@ -28,13 +28,13 @@
namespace caf {
namespace opencl {
template <typename Signature>
template <class... Ts>
class actor_facade;
/// @brief A wrapper for OpenCL's cl_program.
class program {
template <typename Signature>
template <class... Ts>
friend class actor_facade;
public:
......@@ -45,11 +45,13 @@ public:
const char* options = nullptr, uint32_t device_id = 0);
private:
program(context_ptr context, command_queue_ptr queue, program_ptr program);
program(context_ptr context, command_queue_ptr queue, program_ptr program,
std::map<std::string,kernel_ptr> available_kernels);
context_ptr context_;
program_ptr program_;
command_queue_ptr queue_;
std::map<std::string,kernel_ptr> available_kernels_;
};
} // namespace opencl
......
......@@ -23,43 +23,109 @@
#include <algorithm>
#include <functional>
#include "caf/actor_cast.hpp"
#include "caf/optional.hpp"
#include "caf/actor_cast.hpp"
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/actor_facade.hpp"
#include "caf/opencl/spawn_config.hpp"
#include "caf/opencl/opencl_metainfo.hpp"
namespace caf {
namespace detail {
// converts C arrays, i.e., pointers, to vectors
template <typename T>
struct carr_to_vec {
using type = T;
};
struct tuple_construct { };
template <typename T>
struct carr_to_vec<T*> {
using type = std::vector<T>;
template <class... Ts>
struct cl_spawn_helper {
using impl = opencl::actor_facade<Ts...>;
using map_in_fun = std::function<optional<message> (message&)>;
using map_out_fun = typename impl::output_mapping;
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg, Ts&&... xs) const {
return actor_cast<actor>(impl::create(p, fn, cfg,
map_in_fun{}, map_out_fun{},
std::move(xs)...));
}
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg,
map_in_fun map_input, map_out_fun map_output,
Ts&&... xs) const {
return actor_cast<actor>(impl::create(p, fn, cfg, std::move(map_input),
std::move(map_output),
std::move(xs)...));
}
// used by the deprecated spawn_helper
template <class Tuple, long... Is>
actor operator()(tuple_construct,
const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg,
Tuple&& xs,
detail::int_list<Is...>) const {
return actor_cast<actor>(impl::create(p, fn, cfg,
map_in_fun{}, map_out_fun{},
std::move(std::get<Is>(xs))...));
}
template <class Tuple, long... Is>
actor operator()(tuple_construct,
const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg,
map_in_fun map_input, map_out_fun map_output,
Tuple&& xs,
detail::int_list<Is...>) const {
return actor_cast<actor>(impl::create(p, fn, cfg,std::move(map_input),
std::move(map_output),
std::move(std::get<Is>(xs))...));
}
};
template <class Signature>
struct cl_spawn_helper;
template <typename R, typename... Ts>
struct cl_spawn_helper<R(Ts...)> {
using res_t = typename carr_to_vec<R>::type;
using impl = opencl::actor_facade<res_t (typename carr_to_vec<Ts>::type...)>;
using map_arg_fun = std::function<optional<message> (message&)>;
using map_res_fun = typename impl::result_mapping;
template <typename... Us>
actor operator()(const opencl::program& p, const char* fn, Us&&... vs) const {
return actor_cast<actor>(impl::create(p, fn, std::forward<Us>(vs)...));
struct cl_spawn_helper_deprecated;
template <class R, class... Ts>
struct cl_spawn_helper_deprecated<R(Ts...)> {
using input_types
= detail::type_list<typename opencl::carr_to_vec<Ts>::type...>;
using wrapped_input_types
= typename detail::tl_map<input_types,opencl::to_input_arg>::type;
using wrapped_output_type
= opencl::out<typename opencl::carr_to_vec<R>::type>;
using all_types =
typename detail::tl_push_back<
wrapped_input_types,
wrapped_output_type
>::type;
using values = typename detail::tl_apply<all_types, std::tuple>::type;
using helper_type =
typename detail::tl_apply<all_types, cl_spawn_helper>::type;
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& conf,
size_t result_size) const {
values xs;
std::get<tl_size<all_types>::value - 1>(xs) =
wrapped_output_type{[result_size](Ts&...){ return result_size; }};
helper_type f;
return f(tuple_construct{}, p, fn, conf, std::move(xs),
detail::get_indices(xs));
}
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& conf,
size_t result_size,
typename helper_type::map_in_fun map_args,
typename helper_type::map_out_fun map_results) const {
values xs;
std::get<tl_size<all_types>::value - 1>(xs) =
wrapped_output_type{[result_size](Ts&...){ return result_size; }};
auto indices = detail::get_indices(xs);
helper_type f;
return f(tuple_construct{}, p, fn, conf,
std::move(map_args), std::move(map_results),
std::move(xs), std::move(indices));
}
};
......@@ -69,15 +135,131 @@ struct cl_spawn_helper<R(Ts...)> {
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template <class T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::spawn_config& config,
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(prog, fname, config, std::move(x), std::move(xs)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template <class T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn_cl(const char* source,
const char* fname,
const opencl::spawn_config& config,
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(opencl::program::create(source), fname, config,
std::move(x), std::move(xs)...);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template <class Fun, class... Ts>
actor spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args,
Fun map_result,
Ts... xs) {
detail::cl_spawn_helper<Ts...> f;
return f(prog, fname, config, std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template <class Fun, class... Ts>
actor spawn_cl(const char* source,
const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args,
Fun map_result,
Ts... xs) {
detail::cl_spawn_helper<Ts...> f;
return f(opencl::program::create(source), fname, config,
std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
// !!! Below are the deprecated spawn_cl functions !!!
// Signature first to mark them deprcated, gcc will complain otherwise
template <class Signature>
actor spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
detail::cl_spawn_helper<Signature> f;
return f(prog, fname, dims, offset, local_dims, result_size);
size_t result_size = 0) CAF_DEPRECATED;
template <class Signature>
actor spawn_cl(const char* source,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) CAF_DEPRECATED;
template <class Signature, class Fun>
actor spawn_cl(const opencl::program& prog,
const char* fname,
std::function<optional<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) CAF_DEPRECATED;
template <class Signature, class Fun>
actor spawn_cl(const char* source,
const char* fname,
std::function<optional<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) CAF_DEPRECATED;
// now the implementations
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template <class Signature>
actor spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(prog, fname, opencl::spawn_config{dims, offset, local_dims},
result_size);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
......@@ -89,11 +271,12 @@ template <class Signature>
actor spawn_cl(const char* source,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
return spawn_cl<Signature>(opencl::program::create(source), fname,
dims, offset, local_dims, result_size);
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(opencl::program::create(source), fname,
opencl::spawn_config{dims, offset, local_dims}, result_size);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
......@@ -106,12 +289,12 @@ actor spawn_cl(const opencl::program& prog,
std::function<optional<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
detail::cl_spawn_helper<Signature> f;
return f(prog, fname, dims, offset, local_dims, result_size,
std::move(map_args), std::move(map_result));
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(prog, fname, opencl::spawn_config{dims, offset, local_dims},
result_size, std::move(map_args), std::move(map_result));
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
......@@ -125,11 +308,12 @@ actor spawn_cl(const char* source,
std::function<optional<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
detail::cl_spawn_helper<Signature> f;
return f(opencl::program::create(source), fname, dims, offset, local_dims,
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(opencl::program::create(source), fname,
opencl::spawn_config{dims, offset, local_dims},
result_size, std::move(map_args), std::move(map_result));
}
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_SPAWN_CONFIG_HPP
#define CAF_OPENCL_SPAWN_CONFIG_HPP
#include "caf/opencl/global.hpp"
namespace caf {
namespace opencl {
class spawn_config {
public:
spawn_config(const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {})
: dims_{dims},
offset_{offset},
local_dims_{local_dims} {
// nop
}
spawn_config(opencl::dim_vec&& dims,
opencl::dim_vec&& offset = {},
opencl::dim_vec&& local_dims = {})
: dims_{std::move(dims)},
offset_{std::move(offset)},
local_dims_{std::move(local_dims)} {
// nop
}
const opencl::dim_vec& dimensions() const {
return dims_;
}
const opencl::dim_vec& offsets() const {
return offset_;
}
const opencl::dim_vec& local_dimensions() const {
return local_dims_;
}
private:
const opencl::dim_vec dims_;
const opencl::dim_vec offset_;
const opencl::dim_vec local_dims_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_SPAWN_CONFIG_HPP
......@@ -28,6 +28,7 @@
using namespace std;
using namespace caf;
using namespace caf::opencl;
using detail::limited_vector;
......@@ -152,27 +153,28 @@ void multiplier(event_based_actor* self) {
);
};
auto box_res = [] (fvec& result) -> message {
auto box_res = [] (fvec result) -> message {
return make_message(matrix_type{move(result)});
};
// spawn an opencl actor
// 1st arg: source code of one or more opencl kernels
// 2nd arg: name of the kernel to use
auto worker = spawn_cl<fvec (fvec&, fvec&)>(kernel_source, kernel_name,
// 3rd arg: the opencl function operates on vectors,
// this function converts a tuple of two matrices
// to a tuple of vectors; note that this function returns
// an option (an empty results causes the actor to ignore
// the message)
unbox_args,
// 4th arg: converts the ouptut vector back to a matrix that is then
// 3rd arg: the config specifies how many dimensions the kernel uses and how
// many work items are created, creates matrix_size * matrix_size
// global work items in this case
// 4th arg: the opencl function operates on vectors, this function converts
// a tuple of two matrices to a tuple of vectors; note that this
// function returns an option (an empty results causes the actor to
// ignore the message)
// 5th arg: converts the ouptut vector back to a matrix that is then
// used as response message
box_res,
// 5th arg: global dimension arguments for opencl's enqueue,
// creates matrix_size * matrix_size global work items
limited_vector<size_t, 3>{matrix_size, matrix_size}
);
// from 6 : a description of the kernel signature using in/out/in_out classes
// with the argument type packed in vectors
auto worker = spawn_cl(kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
unbox_args, box_res,
in<fvec>{}, in<fvec>{}, out<fvec>{});
// send both matrices to the actor and
// wait for results in form of a matrix_type
......
......@@ -27,11 +27,14 @@
using namespace std;
using namespace caf;
using namespace caf::opencl;
using detail::limited_vector;
namespace {
using fvec = std::vector<float>;
constexpr size_t matrix_size = 8;
constexpr const char* kernel_name = "matrix_mult";
......@@ -55,7 +58,7 @@ constexpr const char* kernel_source = R"__(
} // namespace <anonymous>
void print_as_matrix(const vector<float>& matrix) {
void print_as_matrix(const fvec& matrix) {
for (size_t column = 0; column < matrix_size; ++column) {
for (size_t row = 0; row < matrix_size; ++row) {
cout << fixed << setprecision(2) << setw(9)
......@@ -68,8 +71,8 @@ void print_as_matrix(const vector<float>& matrix) {
void multiplier(event_based_actor* self) {
// the opencl actor only understands vectors
// so these vectors represent the matrices
vector<float> m1(matrix_size * matrix_size);
vector<float> m2(matrix_size * matrix_size);
fvec m1(matrix_size * matrix_size);
fvec m2(matrix_size * matrix_size);
// fill each with ascending values
iota(m1.begin(), m1.end(), 0);
......@@ -81,25 +84,21 @@ void multiplier(event_based_actor* self) {
cout << endl;
// spawn an opencl actor
// template parameter: signature of opencl kernel using proper return type
// instead of output parameter (implicitly
// mapped to the last kernel argument)
// 1st arg: source code of one or more kernels
// 2nd arg: name of the kernel to use
// 3rd arg: global dimension arguments for opencl's enqueue;
// 3rd arg: a spawn configuration that includes:
// - the global dimension arguments for opencl's enqueue
// creates matrix_size * matrix_size global work items
// 4th arg: offsets for global dimensions (optional)
// 5th arg: local dimensions (optional)
// 6th arg: number of elements in the result buffer
auto worker = spawn_cl<float*(float*,float*)>(kernel_source,
kernel_name,
limited_vector<size_t, 3>{matrix_size, matrix_size},
{},
{},
matrix_size * matrix_size);
// - offsets for global dimensions (optional)
// - local dimensions (optional)
// 4th to Nth arg: the kernel signature described by in/out/in_out classes
// that contain the argument type in their template, requires vectors
auto worker = spawn_cl(kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
in<fvec>{}, in<fvec>{}, out<fvec>{});
// send both matrices to the actor and wait for a result
self->sync_send(worker, move(m1), move(m2)).then(
[](const vector<float>& result) {
[](const fvec& result) {
cout << "result: " << endl;
print_as_matrix(result);
}
......@@ -107,7 +106,7 @@ void multiplier(event_based_actor* self) {
}
int main() {
announce<vector<float>>("float_vector");
announce<fvec>("float_vector");
spawn(multiplier);
await_all_actors_done();
shutdown();
......
......@@ -17,6 +17,7 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include <map>
#include <vector>
#include <string>
#include <cstring>
......@@ -34,10 +35,12 @@ namespace caf {
namespace opencl {
program::program(context_ptr context, command_queue_ptr queue,
program_ptr program)
program_ptr program,
map<string, kernel_ptr> available_kernels)
: context_(move(context)),
program_(move(program)),
queue_(move(queue)) {}
queue_(move(queue)),
available_kernels_(move(available_kernels)) {}
program program::create(const char* kernel_source, const char* options,
uint32_t device_id) {
......@@ -91,7 +94,35 @@ program program::create(const char* kernel_source, const char* options,
#endif
throw runtime_error(oss.str());
}
return {context, devices[device_id].cmd_queue_, pptr};
cl_uint number_of_kernels = 0;
clCreateKernelsInProgram(pptr.get(), 0, nullptr, &number_of_kernels);
vector<cl_kernel> kernels(number_of_kernels);
err = clCreateKernelsInProgram(pptr.get(), number_of_kernels, kernels.data(),
nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clCreateKernelsInProgram: " << get_opencl_error(err);
throw runtime_error(oss.str());
}
map<string, kernel_ptr> available_kernels;
for (cl_uint i = 0; i < number_of_kernels; ++i) {
size_t ret_size;
clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, 0, nullptr, &ret_size);
vector<char> name(ret_size);
err = clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, ret_size,
(void*) name.data(), nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetKernelInfo (CL_KERNEL_FUNCTION_NAME): "
<< get_opencl_error(err);
throw runtime_error(oss.str());
}
kernel_ptr kernel;
kernel.reset(move(kernels[i]));
available_kernels.emplace(string(name.data()), move(kernel));
}
return {context, devices[device_id].cmd_queue_, pptr,
move(available_kernels)};
}
} // namespace opencl
......
......@@ -18,16 +18,17 @@ using detail::limited_vector;
namespace {
using ivec = std::vector<int>;
using dims = opencl::dim_vec;
constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32;
constexpr int magic_number = 23;
constexpr size_t problem_size = 1024;
constexpr const char* kernel_name = "matrix_square";
constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_reduce = "reduce";
constexpr const char* kernel_name_const = "const_mod";
constexpr const char* kernel_name_inout = "times_two";
constexpr const char* compiler_flag = "-D CAF_OPENCL_TEST_FLAG";
......@@ -47,7 +48,7 @@ constexpr const char* kernel_source = R"__(
constexpr const char* kernel_source_error = R"__(
__kernel void missing(__global int*) {
size_t semicolon
size_t semicolon_missing
}
)__";
......@@ -94,6 +95,13 @@ constexpr const char* kernel_source_const = R"__(
}
)__";
constexpr const char* kernel_source_inout = R"__(
__kernel void times_two(__global int* values) {
size_t idx = get_global_id(0);
values[idx] = values[idx] * 2;
}
)__";
} // namespace <anonymous>
template<size_t Size>
......@@ -182,11 +190,24 @@ bool operator!=(const square_matrix<Size>& lhs,
using matrix_type = square_matrix<matrix_size>;
size_t get_max_workgroup_size(size_t device_id, size_t dimension) {
size_t max_size = 512;
auto devices = opencl_metainfo::instance()->get_devices()[device_id];
size_t dimsize = devices.get_max_work_items_per_dim()[dimension];
return max_size < dimsize ? max_size : dimsize;
template <class T>
void check_vector_results(const std::string& description,
const std::vector<T>& expected,
const std::vector<T>& result) {
auto cond = (expected == result);
CAF_CHECK(cond);
if (!cond) {
CAF_TEST_INFO(description << " failed.");
std::cout << "Expected: " << std::endl;
for (size_t i = 0; i < expected.size(); ++i) {
std::cout << " " << expected[i];
}
std::cout << std::endl << "Received: " << std::endl;
for (size_t i = 0; i < result.size(); ++i) {
std::cout << " " << result[i];
}
std::cout << std::endl;
}
}
void test_opencl() {
......@@ -195,23 +216,33 @@ void test_opencl() {
152, 174, 196, 218,
248, 286, 324, 362,
344, 398, 452, 506};
auto w1 = spawn_cl<ivec (ivec&)>(program::create(kernel_source),
kernel_name,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
auto w1 = spawn_cl(program::create(kernel_source), kernel_name,
opencl::spawn_config{dims{matrix_size, matrix_size}},
opencl::in<ivec>{}, opencl::out<ivec>{});
self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
self->receive (
[&](const ivec& result) {
CAF_CHECK(result == expected1);
check_vector_results("Simple matrix multiplication using vectors"
"(kernel wrapped in program)",
expected1, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
auto w2 = spawn_cl<ivec (ivec&)>(kernel_source, kernel_name,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
opencl::spawn_config cfg2{dims{matrix_size, matrix_size}};
auto w2 = spawn_cl(kernel_source, kernel_name, cfg2,
opencl::in<ivec>{}, opencl::out<ivec>{});
self->send(w2, make_iota_vector<int>(matrix_size * matrix_size));
self->receive (
[&](const ivec& result) {
CAF_CHECK(result == expected1);
check_vector_results("Simple matrix multiplication using vectors",
expected1, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
const matrix_type expected2(std::move(expected1));
......@@ -222,63 +253,95 @@ void test_opencl() {
}
);
};
auto map_res = [](ivec& result) -> message {
auto map_res = [=](ivec result) -> message {
return make_message(matrix_type{std::move(result)});
};
auto w3 = spawn_cl<ivec (ivec&)>(program::create(kernel_source),
kernel_name, map_arg, map_res,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
opencl::spawn_config cfg3{dims{matrix_size, matrix_size}};
auto w3 = spawn_cl(program::create(kernel_source), kernel_name, cfg3,
map_arg, map_res,
opencl::in<ivec>{}, opencl::out<ivec>{});
self->send(w3, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
CAF_CHECK(expected2 == result);
check_vector_results("Matrix multiplication with user defined type "
"(kernel wrapped in program)",
expected2.data(), result.data());
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
auto w4 = spawn_cl<ivec (ivec&)>(kernel_source, kernel_name,
opencl::spawn_config cfg4{dims{matrix_size, matrix_size}};
auto w4 = spawn_cl(kernel_source, kernel_name, cfg4,
map_arg, map_res,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
opencl::in<ivec>{}, opencl::out<ivec>{});
self->send(w4, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
CAF_CHECK(expected2 == result);
check_vector_results("Matrix multiplication with user defined type",
expected2.data(), result.data());
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
CAF_TEST_INFO("Expecting exception (compiling invalid kernel, "
"semicolon is missing).");
try {
auto create_error = program::create(kernel_source_error);
}
catch (const std::exception& exc) {
CAF_MESSAGE(exc.what());
CAF_CHECK(strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what()) == 0);
auto cond = (strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE",
exc.what()) == 0);
CAF_CHECK(cond);
if (!cond) {
CAF_TEST_INFO("Wrong exception cought for program build failure.");
}
}
// test for opencl compiler flags
auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag);
auto w5 = spawn_cl<ivec (ivec&)>(prog5, kernel_name_compiler_flag,
limited_vector<size_t, 3>{array_size});
opencl::spawn_config cfg5{dims{array_size}};
auto w5 = spawn_cl(prog5, kernel_name_compiler_flag, cfg5,
opencl::in<ivec>{}, opencl::out<ivec>{});
self->send(w5, make_iota_vector<int>(array_size));
auto expected3 = make_iota_vector<int>(array_size);
self->receive (
[&](const ivec& result) {
CAF_CHECK(result == expected3);
check_vector_results("Passing compiler flags", expected3, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
auto get_max_workgroup_size = [](size_t dev_id, size_t dims) -> size_t {
size_t max_size = 512;
auto devices = opencl_metainfo::instance()->get_devices()[dev_id];
size_t dimsize = devices.get_max_work_items_per_dim()[dims];
return max_size < dimsize ? max_size : dimsize;
};
// test for manuel return size selection (max workgroup size 1d)
const int max_workgroup_size = static_cast<int>(get_max_workgroup_size(0,1));
const int reduce_buffer_size = max_workgroup_size * 8;
const int reduce_local_size = max_workgroup_size;
const int reduce_work_groups = reduce_buffer_size / reduce_local_size;
const int reduce_global_size = reduce_buffer_size;
const int reduce_result_size = reduce_work_groups;
ivec arr6(static_cast<size_t>(reduce_buffer_size));
const size_t reduce_buffer_size = static_cast<size_t>(max_workgroup_size) * 8;
const size_t reduce_local_size = static_cast<size_t>(max_workgroup_size);
const size_t reduce_work_groups = reduce_buffer_size / reduce_local_size;
const size_t reduce_global_size = reduce_buffer_size;
const size_t reduce_result_size = reduce_work_groups;
ivec arr6(reduce_buffer_size);
int n = static_cast<int>(arr6.capacity());
std::generate(arr6.begin(), arr6.end(), [&]{ return --n; });
auto w6 = spawn_cl<ivec (ivec&)>(kernel_source_reduce, kernel_name_reduce,
limited_vector<size_t, 3>{static_cast<size_t>(reduce_global_size)},
{},
limited_vector<size_t, 3>{static_cast<size_t>(reduce_local_size)},
static_cast<size_t>(reduce_result_size));
opencl::spawn_config cfg6{dims{reduce_global_size},
dims{ /* no offsets */ },
dims{reduce_local_size}};
auto get_result_size_6 = [=](const ivec&) {
return reduce_result_size;
};
auto w6 = spawn_cl(kernel_source_reduce, kernel_name_reduce, cfg6,
opencl::in<ivec>{}, opencl::out<ivec>{get_result_size_6});
self->send(w6, move(arr6));
ivec expected4{max_workgroup_size * 7, max_workgroup_size * 6,
max_workgroup_size * 5, max_workgroup_size * 4,
......@@ -286,19 +349,51 @@ void test_opencl() {
max_workgroup_size, 0};
self->receive(
[&](const ivec& result) {
CAF_CHECK(result == expected4);
check_vector_results("Passing size for the output", expected4, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
// calculator function for getting the size of the output
auto get_result_size_7 = [=](const ivec&) {
return problem_size;
};
// constant memory arguments
const ivec arr7{magic_number};
auto w7 = spawn_cl<ivec (ivec&)>(kernel_source_const, kernel_name_const,
limited_vector<size_t, 3>{magic_number});
const ivec arr7{problem_size};
auto w7 = spawn_cl(kernel_source_const, kernel_name_const,
opencl::spawn_config{dims{problem_size}},
opencl::in<ivec>{},
opencl::out<ivec>{get_result_size_7});
self->send(w7, move(arr7));
ivec expected5(magic_number);
fill(begin(expected5), end(expected5), magic_number);
ivec expected5(problem_size);
fill(begin(expected5), end(expected5), static_cast<int>(problem_size));
self->receive(
[&](const ivec& result) {
check_vector_results("Using const input argument", expected5, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
// test in_out argument type
ivec input9 = make_iota_vector<int>(problem_size);
ivec expected9{input9};
std::transform(begin(expected9), end(expected9), begin(expected9),
[](const int& val){ return val * 2; });
auto w9 = spawn_cl(kernel_source_inout, kernel_name_inout,
spawn_config{dims{problem_size}},
opencl::in_out<ivec>{});
self->send(w9, std::move(input9));
self->receive(
[&](const ivec& result) {
CAF_CHECK(result == expected5);
check_vector_results("Testing in_out arugment", expected9, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
}
......@@ -310,4 +405,3 @@ CAF_TEST(test_opencl) {
await_all_actors_done();
shutdown();
}
#define CAF_SUITE opencl
#include "caf/test/unit_test.hpp"
#include <vector>
#include <iomanip>
#include <cassert>
#include <iostream>
#include <algorithm>
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
CAF_PUSH_NO_DEPRECATED_WARNING
using namespace caf;
using namespace caf::opencl;
using detail::limited_vector;
namespace {
using ivec = std::vector<int>;
constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32;
constexpr int magic_number = 23;
constexpr const char* kernel_name = "matrix_square";
constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_reduce = "reduce";
constexpr const char* kernel_name_const = "const_mod";
constexpr const char* compiler_flag = "-D CAF_OPENCL_TEST_FLAG";
constexpr const char* kernel_source = R"__(
__kernel void matrix_square(__global int* matrix,
__global int* output) {
size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0);
size_t y = get_global_id(1);
int result = 0;
for (size_t idx = 0; idx < size; ++idx) {
result += matrix[idx + y * size] * matrix[x + idx * size];
}
output[x + y * size] = result;
}
)__";
constexpr const char* kernel_source_error = R"__(
__kernel void missing(__global int*) {
size_t semicolon
}
)__";
constexpr const char* kernel_source_compiler_flag = R"__(
__kernel void compiler_flag(__global int* input,
__global int* output) {
size_t x = get_global_id(0);
# ifdef CAF_OPENCL_TEST_FLAG
output[x] = input[x];
# else
output[x] = 0;
# endif
}
)__";
// http://developer.amd.com/resources/documentation-articles/articles-whitepapers/
// opencl-optimization-case-study-simple-reductions
constexpr const char* kernel_source_reduce = R"__(
__kernel void reduce(__global int* buffer,
__global int* result) {
__local int scratch[512];
int local_index = get_local_id(0);
scratch[local_index] = buffer[get_global_id(0)];
barrier(CLK_LOCAL_MEM_FENCE);
for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
if (local_index < offset) {
int other = scratch[local_index + offset];
int mine = scratch[local_index];
scratch[local_index] = (mine < other) ? mine : other;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if (local_index == 0) {
result[get_group_id(0)] = scratch[0];
}
}
)__";
constexpr const char* kernel_source_const = R"__(
__kernel void const_mod(__constant int* input,
__global int* output) {
size_t idx = get_global_id(0);
output[idx] = input[0];
}
)__";
} // namespace <anonymous>
template<size_t Size>
class square_matrix {
public:
static constexpr size_t num_elements = Size * Size;
static void announce() {
caf::announce<square_matrix>("square_matrix", &square_matrix::data_);
}
square_matrix(square_matrix&&) = default;
square_matrix(const square_matrix&) = default;
square_matrix& operator=(square_matrix&&) = default;
square_matrix& operator=(const square_matrix&) = default;
square_matrix() : data_(num_elements) {
// nop
}
explicit square_matrix(ivec d) : data_(move(d)) {
assert(data_.size() == num_elements);
}
float& operator()(size_t column, size_t row) {
return data_[column + row * Size];
}
const float& operator()(size_t column, size_t row) const {
return data_[column + row * Size];
}
typedef typename ivec::const_iterator const_iterator;
const_iterator begin() const {
return data_.begin();
}
const_iterator end() const {
return data_.end();
}
ivec& data() {
return data_;
}
const ivec& data() const {
return data_;
}
void data(ivec new_data) {
data_ = std::move(new_data);
}
private:
ivec data_;
};
template <class T>
std::vector<T> make_iota_vector(size_t num_elements) {
std::vector<T> result;
result.resize(num_elements);
std::iota(result.begin(), result.end(), T{0});
return result;
}
template <size_t Size>
square_matrix<Size> make_iota_matrix() {
square_matrix<Size> result;
std::iota(result.data().begin(), result.data().end(), 0);
return result;
}
template<size_t Size>
bool operator==(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) {
return lhs.data() == rhs.data();
}
template<size_t Size>
bool operator!=(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) {
return ! (lhs == rhs);
}
using matrix_type = square_matrix<matrix_size>;
template <class T>
void check_vector_results(const std::string& description,
const std::vector<T>& expected,
const std::vector<T>& result) {
auto cond = (expected == result);
CAF_CHECK(cond);
if (!cond) {
CAF_TEST_INFO(description << " failed.");
std::cout << "Expected: " << std::endl;
for (size_t i = 0; i < expected.size(); ++i) {
std::cout << " " << expected[i];
}
std::cout << std::endl << "Received: " << std::endl;
for (size_t i = 0; i < result.size(); ++i) {
std::cout << " " << result[i];
}
std::cout << std::endl;
}
}
void test_opencl_deprecated() {
scoped_actor self;
const ivec expected1{ 56, 62, 68, 74,
152, 174, 196, 218,
248, 286, 324, 362,
344, 398, 452, 506};
auto w1 = spawn_cl<ivec (ivec)>(program::create(kernel_source),
kernel_name,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
self->receive (
[&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors"
"(kernel wrapped in program)",result, expected1);
}
);
auto w2 = spawn_cl<ivec (ivec)>(kernel_source, kernel_name,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w2, make_iota_vector<int>(matrix_size * matrix_size));
self->receive (
[&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors",
expected1, result);
}
);
const matrix_type expected2(std::move(expected1));
auto map_arg = [](message& msg) -> optional<message> {
return msg.apply(
[](matrix_type& mx) {
return make_message(std::move(mx.data()));
}
);
};
auto map_res = [](ivec result) -> message {
return make_message(matrix_type{std::move(result)});
};
auto w3 = spawn_cl<ivec (ivec)>(program::create(kernel_source),
kernel_name, map_arg, map_res,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w3, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type "
"(kernel wrapped in program)",
expected2.data(), result.data());
}
);
auto w4 = spawn_cl<ivec (ivec)>(kernel_source, kernel_name,
map_arg, map_res,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w4, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type",
expected2.data(), result.data());
}
);
CAF_TEST_INFO("Expecting exception (compiling invalid kernel, "
"semicolon is missing).");
try {
auto create_error = program::create(kernel_source_error);
}
catch (const std::exception& exc) {
auto cond = (strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE",
exc.what()) == 0);
CAF_CHECK(cond);
if (!cond) {
CAF_TEST_INFO("Wrong exception cought for program build failure.");
}
}
// test for opencl compiler flags
auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag);
auto w5 = spawn_cl<ivec (ivec)>(prog5, kernel_name_compiler_flag,
limited_vector<size_t, 3>{array_size});
self->send(w5, make_iota_vector<int>(array_size));
auto expected3 = make_iota_vector<int>(array_size);
self->receive (
[&](const ivec& result) {
check_vector_results("Passing compiler flags", expected3, result);
}
);
auto get_max_workgroup_size = [](size_t dev_id, size_t dims) -> size_t {
size_t max_size = 512;
auto devices = opencl_metainfo::instance()->get_devices()[dev_id];
size_t dimsize = devices.get_max_work_items_per_dim()[dims];
return max_size < dimsize ? max_size : dimsize;
};
// test for manuel return size selection (max workgroup size 1d)
const int max_workgroup_size = static_cast<int>(get_max_workgroup_size(0,1));
const int reduce_buffer_size = max_workgroup_size * 8;
const int reduce_local_size = max_workgroup_size;
const int reduce_work_groups = reduce_buffer_size / reduce_local_size;
const int reduce_global_size = reduce_buffer_size;
const int reduce_result_size = reduce_work_groups;
ivec arr6(static_cast<size_t>(reduce_buffer_size));
int n = static_cast<int>(arr6.capacity());
std::generate(arr6.begin(), arr6.end(), [&]{ return --n; });
auto w6 = spawn_cl<ivec (ivec)>(kernel_source_reduce, kernel_name_reduce,
limited_vector<size_t, 3>{static_cast<size_t>(reduce_global_size)},
{},
limited_vector<size_t, 3>{static_cast<size_t>(reduce_local_size)},
static_cast<size_t>(reduce_result_size));
self->send(w6, move(arr6));
ivec expected4{max_workgroup_size * 7, max_workgroup_size * 6,
max_workgroup_size * 5, max_workgroup_size * 4,
max_workgroup_size * 3, max_workgroup_size * 2,
max_workgroup_size, 0};
self->receive(
[&](const ivec& result) {
check_vector_results("Passing size for the output", expected4, result);
}
);
// constant memory arguments
const ivec arr7{magic_number};
auto w7 = spawn_cl<ivec (ivec)>(kernel_source_const, kernel_name_const,
limited_vector<size_t, 3>{magic_number});
self->send(w7, move(arr7));
ivec expected5(magic_number);
fill(begin(expected5), end(expected5), magic_number);
self->receive(
[&](const ivec& result) {
check_vector_results("Using const input argument", expected5, result);
}
);
}
CAF_TEST(test_opencl_deprecated) {
std::cout << "Staring deprecated OpenCL test" << std::endl;
announce<ivec>("ivec");
matrix_type::announce();
test_opencl_deprecated();
await_all_actors_done();
shutdown();
std::cout << "Done with depreacted OpenCL test" << std::endl;
}
CAF_POP_WARNINGS
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