Commit 522e26e3 authored by Joseph Noir's avatar Joseph Noir

Enable kernel executions with multiple stages

This commit introduces memory references to memory buffers on OpenCL
devices and thus allows passing buffers between OpenCL actors on the
same host without the need to transfer memory between GPU and CPU
context. For this purpose, the existing wrappers that tag arguments
as input, output, or input as well as output now accept template
parameters to specify whether the argument is expected and / or
returned as a value (i.e., std::vector) or a reference.

On device computations that only return references reply with a message
containing the references directly and do not wait for the computation
to finish. OpenCL actors that receive such references enqueue their
kernels using OpenCL events to express the dependencies.

The function that maps incoming arguments of OpenCL actors now accepts
an optional additional argument, to allow adaption of the index space
for the execution depending on the arguments.
parent a31c0fc8
...@@ -3,8 +3,7 @@ project(caf_opencl C CXX) ...@@ -3,8 +3,7 @@ project(caf_opencl C CXX)
# get header files; only needed by CMake generators, # get header files; only needed by CMake generators,
# e.g., for creating proper Xcode projects # e.g., for creating proper Xcode projects
file(GLOB LIBCAF_OPENCL_HDRS "caf/opencl/*.hpp") file(GLOB_RECURSE LIBCAF_OPENCL_HDRS "caf/*.hpp")
file(GLOB LIBCAF_OPENCL_HDRS "caf/opencl/detail/*.hpp")
add_custom_target(libcaf_opencl) add_custom_target(libcaf_opencl)
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* 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_ACTOR_FACADE_HPP
#define CAF_OPENCL_ACTOR_FACADE_HPP
#include <ostream>
#include <iostream>
#include <algorithm>
#include <stdexcept>
#include "caf/all.hpp"
#include "caf/intrusive_ptr.hpp"
#include "caf/detail/int_list.hpp"
#include "caf/detail/type_list.hpp"
#include "caf/detail/limited_vector.hpp"
#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 manager;
template <class List>
struct function_sig_from_outputs;
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;
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 monitorable_actor {
public:
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 size_vec = std::vector<size_t>;
using command_type =
typename command_sig_from_outputs<actor_facade, output_types>::type;
const char* name() const override {
return "OpenCL actor";
}
static actor create(actor_config actor_cfg, const program& prog,
const char* kernel_name, const spawn_config& spawn_cfg,
input_mapping map_args, output_mapping map_result,
Ts&&... xs) {
if (spawn_cfg.dimensions().empty()) {
auto str = "OpenCL kernel needs at least 1 global dimension.";
CAF_LOG_ERROR(str);
throw std::runtime_error(str);
}
auto check_vec = [&](const dim_vec& vec, const char* name) {
if (! vec.empty() && vec.size() != spawn_cfg.dimensions().size()) {
std::ostringstream oss;
oss << name << " vector is not empty, but "
<< "its size differs from global dimensions vector's size";
CAF_LOG_ERROR(CAF_ARG(oss.str()));
throw std::runtime_error(oss.str());
}
};
check_vec(spawn_cfg.offsets(), "offsets");
check_vec(spawn_cfg.local_dimensions(), "local dimensions");
auto& sys = actor_cfg.host->system();
auto itr = prog.available_kernels_.find(kernel_name);
if (itr == prog.available_kernels_.end()) {
kernel_ptr kernel;
kernel.reset(v2get(CAF_CLF(clCreateKernel), prog.program_.get(),
kernel_name),
false);
return make_actor<actor_facade, actor>(sys.next_actor_id(), sys.node(),
&sys, std::move(actor_cfg),
prog, kernel, spawn_cfg,
std::move(map_args),
std::move(map_result),
std::forward_as_tuple(xs...));
}
return make_actor<actor_facade, actor>(sys.next_actor_id(), sys.node(),
&sys, std::move(actor_cfg),
prog, itr->second, spawn_cfg,
std::move(map_args),
std::move(map_result),
std::forward_as_tuple(xs...));
}
void enqueue(strong_actor_ptr sender, message_id mid, message content,
execution_unit*) override {
CAF_PUSH_AID(id());
CAF_LOG_TRACE("");
if (map_args_) {
auto mapped = map_args_(content);
if (! mapped) {
return;
}
content = std::move(*mapped);
}
if (! content.match_elements(input_types{})) {
return;
}
auto hdl = std::make_tuple(sender, mid.response_id());
evnt_vec events;
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>(std::move(hdl),
actor_cast<strong_actor_ptr>(this),
std::move(events),
std::move(input_buffers),
std::move(output_buffers),
std::move(result_sizes),
std::move(content));
cmd->enqueue();
}
void enqueue(mailbox_element_ptr ptr, execution_unit* eu) override {
CAF_ASSERT(ptr != nullptr);
CAF_LOG_TRACE(CAF_ARG(*ptr));
enqueue(ptr->sender, ptr->mid, ptr->move_content_to_message(), eu);
}
actor_facade(actor_config actor_cfg,
const program& prog, kernel_ptr kernel,
spawn_config spawn_cfg,
input_mapping map_args, output_mapping map_result,
std::tuple<Ts...> xs)
: monitorable_actor(actor_cfg),
kernel_(std::move(kernel)),
program_(prog.program_),
context_(prog.context_),
queue_(prog.queue_),
spawn_cfg_(std::move(spawn_cfg)),
map_args_(std::move(map_args)),
map_results_(std::move(map_result)),
argument_types_(std::move(xs)) {
CAF_LOG_TRACE(CAF_ARG(this->id()));
default_output_size_ = std::accumulate(spawn_cfg_.dimensions().begin(),
spawn_cfg_.dimensions().end(),
size_t{1},
std::multiplies<size_t>{});
}
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_kernel_arguments(evnt_vec& events, args_vec& input_buffers,
args_vec& output_buffers, size_vec& sizes,
message& msg, detail::int_list<I, Is...>) {
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_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(buffer, false);
input_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<void*>(&input_buffers.back()));
}
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(buffer, false);
output_buffers.push_back(tmp);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(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(), static_cast<unsigned>(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_;
spawn_config spawn_cfg_;
input_mapping map_args_;
output_mapping map_results_;
std::tuple<Ts...> argument_types_;
size_t default_output_size_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_ACTOR_FACADE_HPP
...@@ -27,9 +27,24 @@ ...@@ -27,9 +27,24 @@
#include "caf/message.hpp" #include "caf/message.hpp"
#include "caf/optional.hpp" #include "caf/optional.hpp"
#include "caf/opencl/mem_ref.hpp"
namespace caf { namespace caf {
namespace opencl { namespace opencl {
// Tag classes to mark arguments received in a messages as reference or value
/// Arguments tagged as `val` are expected as a vector (or value in case
/// of a private argument).
struct val { };
/// Arguments tagged as `mref` are expected as mem_ref, which is can be returned
/// by other opencl actors.
struct mref { };
/// Arguments tagged as `hidden` are created by the actor, using the config
/// passed in the argument wrapper. Only available for local and priv arguments.
struct hidden { };
/// Use as a default way to calculate output size. 0 will be set to the number /// Use as a default way to calculate output size. 0 will be set to the number
/// of work items at runtime. /// of work items at runtime.
struct dummy_size_calculator { struct dummy_size_calculator {
...@@ -40,19 +55,36 @@ struct dummy_size_calculator { ...@@ -40,19 +55,36 @@ struct dummy_size_calculator {
}; };
/// Mark an a spawn_cl template argument as input only /// Mark an a spawn_cl template argument as input only
template <class Arg> template <class Arg, class Tag = val>
struct in { struct in {
static_assert(std::is_same<Tag,val>::value || std::is_same<Tag,mref>::value,
"Argument of type `in` must be passed as value or mem_ref.");
using tag_type = Tag;
using arg_type = typename std::decay<Arg>::type; using arg_type = typename std::decay<Arg>::type;
}; };
/// Mark an a spawn_cl template argument as input and output /// Mark an a spawn_cl template argument as input and output
template <class Arg> template <class Arg, class TagIn = val, class TagOut = val>
struct in_out { struct in_out {
static_assert(
std::is_same<TagIn,val>::value || std::is_same<TagIn,mref>::value,
"Argument of type `in_out` must be passed as value or mem_ref."
);
static_assert(
std::is_same<TagOut,val>::value || std::is_same<TagOut,mref>::value,
"Argument of type `in_out` must be returned as value or mem_ref."
);
using tag_in_type = TagIn;
using tag_out_type = TagOut;
using arg_type = typename std::decay<Arg>::type; using arg_type = typename std::decay<Arg>::type;
}; };
template <class Arg> template <class Arg, class Tag = val>
struct out { struct out {
static_assert(std::is_same<Tag,val>::value || std::is_same<Tag,mref>::value,
"Argument of type `out` must be returned as value or mem_ref.");
using tag_type = Tag;
using arg_type = typename std::decay<Arg>::type;
out() = default; out() = default;
template <class F> template <class F>
out(F fun) { out(F fun) {
...@@ -72,6 +104,100 @@ struct out { ...@@ -72,6 +104,100 @@ struct out {
std::function<optional<size_t> (message&)> fun_; std::function<optional<size_t> (message&)> fun_;
}; };
template <class Arg>
struct scratch {
using arg_type = typename std::decay<Arg>::type;
scratch() = default;
template <class F>
scratch(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) : 0UL;
}
std::function<optional<size_t> (message&)> fun_;
};
/// Argument placed in local memory. Cannot be initalized from the CPU, but
/// requires a size that is calculated depending on the input.
template <class Arg>
struct local {
using arg_type = typename std::decay<Arg>::type;
local() = default;
template <class F>
local(size_t size, F fun) : size_(size) {
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;
};
}
local(size_t size) : size_(size) { }
size_t operator()(message& msg) const {
if (fun_) {
auto res = fun_(msg);
if (res)
return *res;
}
return size_;
}
size_t size_;
std::function<optional<size_t> (message&)> fun_;
};
/// Argument placed in private memory. Requires a default value but can
/// alternatively be calculated depending on the input through a function
/// passed to the constructor.
template <class Arg, class Tag = hidden>
struct priv {
static_assert(std::is_same<Tag,val>::value || std::is_same<Tag,hidden>::value,
"Argument of type `priv` must be returned as value or hidden.");
using tag_type = Tag;
using arg_type = typename std::decay<Arg>::type;
priv() = default;
template <class F>
priv(Arg val, F fun) : value_(val) {
static_assert(std::is_same<Tag,hidden>::value,
"Argument of type `priv` can only be initialized with a value"
"if it is manged by the actor, i.e., tagged as hidden.");
fun_ = [fun](message& msg) -> optional<Arg> {
auto res = msg.apply(fun);
Arg result;
if (res) {
res->apply([&](Arg x) { result = x; });
return result;
}
return none;
};
}
priv(Arg val) : value_(val) {
static_assert(std::is_same<Tag,hidden>::value,
"Argument of type `priv` can only be initialized with a value"
"if it is manged by the actor, i.e., tagged as hidden.");
}
Arg operator()(message& msg) const {
if (fun_) {
auto res = fun_(msg);
if (res)
return *res;
}
return value_;
}
Arg value_;
std::function<optional<Arg> (message&)> fun_;
};
///Cconverts C arrays, i.e., pointers, to vectors. ///Cconverts C arrays, i.e., pointers, to vectors.
template <class T> template <class T>
...@@ -88,62 +214,197 @@ struct carr_to_vec<T*> { ...@@ -88,62 +214,197 @@ struct carr_to_vec<T*> {
template <class T> template <class T>
struct is_opencl_arg : std::false_type {}; struct is_opencl_arg : std::false_type {};
template <class T> template <class T, class Tag>
struct is_opencl_arg<in<T>> : std::true_type {}; struct is_opencl_arg<in<T, Tag>> : std::true_type {};
template <class T, class TagIn, class TagOut>
struct is_opencl_arg<in_out<T, TagIn, TagOut>> : std::true_type {};
template <class T, class Tag>
struct is_opencl_arg<out<T, Tag>> : std::true_type {};
template <class T> template <class T>
struct is_opencl_arg<in_out<T>> : std::true_type {}; struct is_opencl_arg<scratch<T>> : std::true_type {};
template <class T> template <class T>
struct is_opencl_arg<out<T>> : std::true_type {}; struct is_opencl_arg<local<T>> : std::true_type {};
template <class T, class Tag>
struct is_opencl_arg<priv<T, Tag>> : std::true_type {};
/// Filter type lists for input arguments, in and in_out. /// Filter type lists for input arguments
template <class T> template <class T>
struct is_input_arg : std::false_type {}; struct is_input_arg : std::false_type {};
template <class T> template <class T, class Tag>
struct is_input_arg<in<T>> : std::true_type {}; struct is_input_arg<in<T, Tag>> : std::true_type {};
template <class T, class TagIn, class TagOut>
struct is_input_arg<in_out<T, TagIn, TagOut>> : std::true_type {};
template <class T> template <class T>
struct is_input_arg<in_out<T>> : std::true_type {}; struct is_input_arg<priv<T,val>> : std::true_type {};
/// Filter type lists for output arguments, in_out and out. /// Filter type lists for output arguments
template <class T> template <class T>
struct is_output_arg : std::false_type {}; struct is_output_arg : std::false_type {};
template <class T, class Tag>
struct is_output_arg<out<T, Tag>> : std::true_type {};
template <class T, class TagIn, class TagOut>
struct is_output_arg<in_out<T, TagIn, TagOut>> : std::true_type {};
/// Filter for arguments that require size
template <class T> template <class T>
struct is_output_arg<out<T>> : std::true_type {}; struct requires_size_arg : std::false_type {};
template <class T, class Tag>
struct requires_size_arg<out<T, Tag>> : std::true_type {};
template <class T> template <class T>
struct is_output_arg<in_out<T>> : std::true_type {}; struct requires_size_arg<scratch<T>> : std::true_type {};
/// Filter for arguments that require size, in this case only out.
template <class T> template <class T>
struct requires_size_arg : std::false_type {}; struct requires_size_arg<local<T>> : std::true_type {};
//template <class T, class Tag>
//struct requires_size_arg<priv<T, Tag>> : std::true_type {};
/// Filter mem_refs
template <class T> template <class T>
struct requires_size_arg<out<T>> : std::true_type {}; struct is_ref_type : std::false_type {};
template <class T>
struct is_ref_type<mem_ref<T>> : std::true_type {};
template <class T>
struct is_val_type : std::true_type {};
template <class T>
struct is_val_type<mem_ref<T>> : std::false_type {};
/// extract types /// extract types
template <class T> template <class T>
struct extract_type { }; struct extract_type { };
template <class T> template <class T, class Tag>
struct extract_type<in<T>> { struct extract_type<in<T, Tag>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type;
};
template <class T, class TagIn, class TagOut>
struct extract_type<in_out<T, TagIn, TagOut>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type;
};
template <class T, class Tag>
struct extract_type<out<T, Tag>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type; using type = typename std::decay<typename carr_to_vec<T>::type>::type;
}; };
template <class T> template <class T>
struct extract_type<in_out<T>> { struct extract_type<scratch<T>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type; using type = typename std::decay<typename carr_to_vec<T>::type>::type;
}; };
template <class T> template <class T>
struct extract_type<out<T>> { struct extract_type<local<T>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type;
};
template <class T, class Tag>
struct extract_type<priv<T, Tag>> {
using type = typename std::decay<typename carr_to_vec<T>::type>::type; using type = typename std::decay<typename carr_to_vec<T>::type>::type;
}; };
/// extract type expected in an incoming message
template <class T>
struct extract_input_type { };
template <class Arg>
struct extract_input_type<in<Arg, val>> {
using type = std::vector<Arg>;
};
template <class Arg>
struct extract_input_type<in<Arg, mref>> {
using type = opencl::mem_ref<Arg>;
};
template <class Arg, class TagOut>
struct extract_input_type<in_out<Arg, val, TagOut>> {
using type = std::vector<Arg>;
};
template <class Arg, class TagOut>
struct extract_input_type<in_out<Arg, mref, TagOut>> {
using type = opencl::mem_ref<Arg>;
};
template <class Arg>
struct extract_input_type<priv<Arg,val>> {
using type = Arg;
};
/// extract type sent in an outgoing message
template <class T>
struct extract_output_type { };
template <class Arg>
struct extract_output_type<out<Arg, val>> {
using type = std::vector<Arg>;
};
template <class Arg>
struct extract_output_type<out<Arg, mref>> {
using type = opencl::mem_ref<Arg>;
};
template <class Arg, class TagIn>
struct extract_output_type<in_out<Arg, TagIn, val>> {
using type = std::vector<Arg>;
};
template <class Arg, class TagIn>
struct extract_output_type<in_out<Arg, TagIn, mref>> {
using type = opencl::mem_ref<Arg>;
};
/// extract input tag
template <class T>
struct extract_input_tag { };
template <class Arg, class Tag>
struct extract_input_tag<in<Arg, Tag>> {
using tag = Tag;
};
template <class Arg, class TagIn, class TagOut>
struct extract_input_tag<in_out<Arg, TagIn, TagOut>> {
using tag = TagIn;
};
template <class Arg>
struct extract_input_tag<priv<Arg,val>> {
using tag = val;
};
/// extract output tag
template <class T>
struct extract_output_tag { };
template <class Arg, class Tag>
struct extract_output_tag<out<Arg, Tag>> {
using tag = Tag;
};
template <class Arg, class TagIn, class TagOut>
struct extract_output_tag<in_out<Arg, TagIn, TagOut>> {
using tag = TagOut;
};
/// Create the return message from tuple arumgent /// Create the return message from tuple arumgent
struct message_from_results { struct message_from_results {
template <class T, class... Ts> template <class T, class... Ts>
...@@ -156,6 +417,106 @@ struct message_from_results { ...@@ -156,6 +417,106 @@ struct message_from_results {
} }
}; };
/// Calculate output indices from the kernel message
// index in output tuple
template <int Counter, class Arg>
struct out_index_of {
static constexpr int value = -1;
static constexpr int next = Counter;
};
template <int Counter, class Arg, class TagIn, class TagOut>
struct out_index_of<Counter, in_out<Arg,TagIn,TagOut>> {
static constexpr int value = Counter;
static constexpr int next = Counter + 1;
};
template <int Counter, class Arg, class Tag>
struct out_index_of<Counter, out<Arg,Tag>> {
static constexpr int value = Counter;
static constexpr int next = Counter + 1;
};
// index in input message
template <int Counter, class Arg>
struct in_index_of {
static constexpr int value = -1;
static constexpr int next = Counter;
};
template <int Counter, class Arg, class Tag>
struct in_index_of<Counter, in<Arg,Tag>> {
static constexpr int value = Counter;
static constexpr int next = Counter + 1;
};
template <int Counter, class Arg, class TagIn, class TagOut>
struct in_index_of<Counter, in_out<Arg,TagIn,TagOut>> {
static constexpr int value = Counter;
static constexpr int next = Counter + 1;
};
template <int Counter, class Arg>
struct in_index_of<Counter, priv<Arg,val>> {
static constexpr int value = Counter;
static constexpr int next = Counter + 1;
};
template <int In, int Out, class T>
struct cl_arg_info {
static constexpr int in_pos = In;
static constexpr int out_pos = Out;
using type = T;
};
template <class ListA, class ListB, int InCounter, int OutCounter>
struct cl_arg_info_list_impl;
template <class Arg, class... Remaining, int InCounter, int OutCounter>
struct cl_arg_info_list_impl<detail::type_list<>,
detail::type_list<Arg, Remaining...>,
InCounter, OutCounter> {
using in_idx = in_index_of<InCounter, Arg>;
using out_idx = out_index_of<OutCounter, Arg>;
using type =
typename cl_arg_info_list_impl<
detail::type_list<cl_arg_info<in_idx::value, out_idx::value, Arg>>,
detail::type_list<Remaining...>,
in_idx::next, out_idx::next
>::type;
};
template <class... Args, class Arg, class... Remaining,
int InCounter, int OutCounter>
struct cl_arg_info_list_impl<detail::type_list<Args...>,
detail::type_list<Arg, Remaining...>,
InCounter, OutCounter> {
using in_idx = in_index_of<InCounter, Arg>;
using out_idx = out_index_of<OutCounter, Arg>;
using type =
typename cl_arg_info_list_impl<
detail::type_list<Args..., cl_arg_info<in_idx::value, out_idx::value, Arg>>,
detail::type_list<Remaining...>,
in_idx::next, out_idx::next
>::type;
};
template <class... Args, int InCounter, int OutCounter>
struct cl_arg_info_list_impl<detail::type_list<Args...>, detail::type_list<>,
InCounter, OutCounter> {
using type = detail::type_list<Args...>;
};
template <class List>
struct cl_arg_info_list {
using type = typename cl_arg_info_list_impl<
detail::type_list<>,
List, 0, 0
>::type;
};
/// Helpers for conversion in deprecated spawn functions /// Helpers for conversion in deprecated spawn functions
template <class T> template <class T>
......
...@@ -37,157 +37,244 @@ ...@@ -37,157 +37,244 @@
#include "caf/opencl/arguments.hpp" #include "caf/opencl/arguments.hpp"
#include "caf/opencl/smart_ptr.hpp" #include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp" #include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/spawn_config.hpp"
namespace caf { namespace caf {
namespace opencl { namespace opencl {
template <class FacadeType, class... Ts> /// A command represents the execution of a kernel on a device. It handles the
/// OpenCL calls to enqueue the kernel with the index space and keeps references
/// to the management data during the execution. Furthermore, the command sends
/// the execution results to the responsible actor.
template <class Actor, class... Ts>
class command : public ref_counted { class command : public ref_counted {
public: public:
command(std::tuple<strong_actor_ptr,message_id> handle, using result_types = detail::type_list<Ts...>;
strong_actor_ptr actor_facade,
std::vector<cl_event> events, std::vector<mem_ptr> input_buffers, command(response_promise promise,
std::vector<mem_ptr> output_buffers, std::vector<size_t> result_sizes, strong_actor_ptr parent,
message msg) std::vector<cl_event> events,
: result_sizes_(std::move(result_sizes)), std::vector<cl_mem_ptr> inputs,
handle_(std::move(handle)), std::vector<cl_mem_ptr> outputs,
actor_facade_(std::move(actor_facade)), std::vector<cl_mem_ptr> scratches,
std::vector<size_t> lengths,
message msg,
std::tuple<Ts...> output_tuple,
spawn_config config)
: lengths_(std::move(lengths)),
promise_(std::move(promise)),
cl_actor_(std::move(parent)),
mem_in_events_(std::move(events)), mem_in_events_(std::move(events)),
input_buffers_(std::move(input_buffers)), input_buffers_(std::move(inputs)),
output_buffers_(std::move(output_buffers)), output_buffers_(std::move(outputs)),
msg_(std::move(msg)) { scratch_buffers_(std::move(scratches)),
results_(output_tuple),
msg_(std::move(msg)),
config_(std::move(config)) {
// nop // nop
} }
~command() override { ~command() override {
for (auto& e : mem_in_events_) { for (auto& e : mem_in_events_) {
v1callcl(CAF_CLF(clReleaseEvent),e); if (e)
v1callcl(CAF_CLF(clReleaseEvent), e);
} }
for (auto& e : mem_out_events_) { for (auto& e : mem_out_events_) {
v1callcl(CAF_CLF(clReleaseEvent),e); if (e)
v1callcl(CAF_CLF(clReleaseEvent), e);
} }
} }
void enqueue() { /// Enqueue the kernel for execution, schedule reading of the results and
/// set a callback to send the results to the actor identified by the handle.
/// Only called if the results includes at least one type that is not a
/// mem_ref.
template <class Q = result_types>
typename std::enable_if<
!detail::tl_forall<Q, is_ref_type>::value,
void
>::type
enqueue() {
// Errors in this function can not be handled by opencl_err.hpp // Errors in this function can not be handled by opencl_err.hpp
// because they require non-standard error handling // because they require non-standard error handling
CAF_LOG_TRACE("command::enqueue()"); CAF_LOG_TRACE("command::enqueue() mixed");
this->ref(); // reference held by the OpenCL comand queue this->ref(); // reference held by the OpenCL comand queue
cl_event event_k;
auto data_or_nullptr = [](const dim_vec& vec) { auto data_or_nullptr = [](const dim_vec& vec) {
return vec.empty() ? nullptr : vec.data(); return vec.empty() ? nullptr : vec.data();
}; };
auto actor_facade = auto parent = static_cast<Actor*>(actor_cast<abstract_actor*>(cl_actor_));
static_cast<FacadeType*>(actor_cast<abstract_actor*>(actor_facade_));
// OpenCL expects cl_uint (unsigned int), hence the cast // OpenCL expects cl_uint (unsigned int), hence the cast
mem_out_events_.emplace_back();
cl_int err = clEnqueueNDRangeKernel( cl_int err = clEnqueueNDRangeKernel(
actor_facade->queue_.get(), actor_facade->kernel_.get(), parent->queue_.get(), parent->kernel_.get(),
static_cast<cl_uint>(actor_facade->spawn_cfg_.dimensions().size()), static_cast<unsigned int>(config_.dimensions().size()),
data_or_nullptr(actor_facade->spawn_cfg_.offsets()), data_or_nullptr(config_.offsets()),
data_or_nullptr(actor_facade->spawn_cfg_.dimensions()), data_or_nullptr(config_.dimensions()),
data_or_nullptr(actor_facade->spawn_cfg_.local_dimensions()), data_or_nullptr(config_.local_dimensions()),
static_cast<cl_uint>(mem_in_events_.size()), static_cast<unsigned int>(mem_in_events_.size()),
(mem_in_events_.empty() ? nullptr : mem_in_events_.data()), &event_k (mem_in_events_.empty() ? nullptr : mem_in_events_.data()),
&mem_out_events_.back()
); );
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
CAF_LOG_ERROR("clEnqueueNDRangeKernel: " CAF_LOG_ERROR("clEnqueueNDRangeKernel: "
<< CAF_ARG(get_opencl_error(err))); << CAF_ARG(get_opencl_error(err)));
clReleaseEvent(event_k);
this->deref(); this->deref();
return; return;
} }
enqueue_read_buffers(event_k, detail::get_indices(result_buffers_)); size_t pos = 0;
cl_event marker; CAF_ASSERT(!mem_out_events_.empty());
enqueue_read_buffers(pos, mem_out_events_, detail::get_indices(results_));
CAF_ASSERT(mem_out_events_.size() > 1);
cl_event marker_event;
#if defined(__APPLE__) #if defined(__APPLE__)
err = clEnqueueMarkerWithWaitList( err = clEnqueueMarkerWithWaitList(
actor_facade->queue_.get(), parent->queue_.get(),
static_cast<cl_uint>(mem_out_events_.size()), static_cast<unsigned int>(mem_out_events_.size()),
mem_out_events_.data(), &marker mem_out_events_.data(),
); &marker_event
);
std::string name = "clEnqueueMarkerWithWaitList";
#else #else
err = clEnqueueMarker(actor_facade->queue_.get(), &marker); err = clEnqueueMarker(parent->queue_.get(), &marker_event);
std::string name = "clEnqueueMarker";
#endif #endif
if (err != CL_SUCCESS) { callback_.reset(marker_event, false);
CAF_LOG_ERROR("clSetEventCallback: " << CAF_ARG(get_opencl_error(err))); if (err != CL_SUCCESS) {
clReleaseEvent(marker); CAF_LOG_ERROR(name << ": " << CAF_ARG(get_opencl_error(err)));
clReleaseEvent(event_k); this->deref(); // callback is not set
this->deref(); // callback is not set return;
return; }
} err = clSetEventCallback(callback_.get(), CL_COMPLETE,
err = clSetEventCallback(marker, CL_COMPLETE, [](cl_event, cl_int, void* data) {
[](cl_event, cl_int, void* data) { auto cmd = reinterpret_cast<command*>(data);
auto cmd = reinterpret_cast<command*>(data); cmd->handle_results();
cmd->handle_results(); cmd->deref();
cmd->deref(); },
}, this);
this); if (err != CL_SUCCESS) {
if (err != CL_SUCCESS) { CAF_LOG_ERROR("clSetEventCallback: " << CAF_ARG(get_opencl_error(err)));
CAF_LOG_ERROR("clSetEventCallback: " << CAF_ARG(get_opencl_error(err))); this->deref(); // callback is not set
clReleaseEvent(marker); return;
clReleaseEvent(event_k); }
this->deref(); // callback is not set err = clFlush(parent->queue_.get());
return; if (err != CL_SUCCESS) {
} CAF_LOG_ERROR("clFlush: " << CAF_ARG(get_opencl_error(err)));
err = clFlush(actor_facade->queue_.get()); }
if (err != CL_SUCCESS) {
CAF_LOG_ERROR("clFlush: " << CAF_ARG(get_opencl_error(err)));
}
mem_out_events_.push_back(std::move(event_k));
mem_out_events_.push_back(std::move(marker));
} }
private: /// Enqueue the kernel for execution and send the mem_refs relating to the
std::vector<size_t> result_sizes_; /// results to the next actor. A callback is set to clean up the commmand
std::tuple<strong_actor_ptr,message_id> handle_; /// once the execution is finished. Only called if the results only consist
strong_actor_ptr actor_facade_; /// of mem_ref types.
std::vector<cl_event> mem_in_events_; template <class Q = result_types>
std::vector<cl_event> mem_out_events_; typename std::enable_if<
std::vector<mem_ptr> input_buffers_; detail::tl_forall<Q, is_ref_type>::value,
std::vector<mem_ptr> output_buffers_; void
std::tuple<Ts...> result_buffers_; >::type
message msg_; // required to keep the argument buffers alive (async copy) enqueue() {
// Errors in this function can not be handled by opencl_err.hpp
void enqueue_read_buffers(cl_event&, detail::int_list<>) { // because they require non-standard error handling
// nop, end of recursion CAF_LOG_TRACE("command::enqueue() all references");
this->ref(); // reference held by the OpenCL command queue
auto data_or_nullptr = [](const dim_vec& vec) {
return vec.empty() ? nullptr : vec.data();
};
auto parent = static_cast<Actor*>(actor_cast<abstract_actor*>(cl_actor_));
cl_event execution_event;
cl_int err = clEnqueueNDRangeKernel(
parent->queue_.get(), parent->kernel_.get(),
static_cast<cl_uint>(config_.dimensions().size()),
data_or_nullptr(config_.offsets()),
data_or_nullptr(config_.dimensions()),
data_or_nullptr(config_.local_dimensions()),
static_cast<unsigned int>(mem_in_events_.size()),
(mem_in_events_.empty() ? nullptr : mem_in_events_.data()),
&execution_event
);
callback_.reset(execution_event, false);
if (err != CL_SUCCESS) {
CAF_LOG_ERROR("clEnqueueNDRangeKernel: "
<< CAF_ARG(get_opencl_error(err)));
this->deref();
return;
}
err = clSetEventCallback(callback_.get(), CL_COMPLETE,
[](cl_event, cl_int, void* data) {
auto c = reinterpret_cast<command*>(data);
c->deref();
},
this);
if (err != CL_SUCCESS) {
CAF_LOG_ERROR("clSetEventCallback: " << CAF_ARG(get_opencl_error(err)));
this->deref(); // callback is not set
return;
}
err = clFlush(parent->queue_.get());
if (err != CL_SUCCESS)
CAF_LOG_ERROR("clFlush: " << CAF_ARG(get_opencl_error(err)));
auto msg = msg_adding_event{callback_}(results_);
promise_.deliver(std::move(msg));
} }
template <long I, long... Is> private:
void enqueue_read_buffers(cl_event& kernel_done, detail::int_list<I, Is...>) { template <long I, class T>
auto actor_facade = void enqueue_read(std::vector<T>&, std::vector<cl_event>& events,
static_cast<FacadeType*>(actor_cast<abstract_actor*>(actor_facade_)); size_t& pos) {
using container_type = auto p = static_cast<Actor*>(actor_cast<abstract_actor*>(cl_actor_));
typename std::tuple_element<I, std::tuple<Ts...>>::type; events.emplace_back();
using value_type = typename container_type::value_type; auto size = lengths_[pos];
cl_event event; auto buffer_size = sizeof(T) * size;
auto size = result_sizes_[I]; std::get<I>(results_).resize(size);
auto buffer_size = sizeof(value_type) * result_sizes_[I]; auto err = clEnqueueReadBuffer(p->queue_.get(), output_buffers_[pos].get(),
std::get<I>(result_buffers_).resize(size);
auto err = clEnqueueReadBuffer(actor_facade->queue_.get(),
output_buffers_[I].get(),
CL_FALSE, 0, buffer_size, CL_FALSE, 0, buffer_size,
std::get<I>(result_buffers_).data(), std::get<I>(results_).data(), 1,
1, &kernel_done, &event); events.data(), &events.back());
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
this->deref(); // failed to enqueue command this->deref(); // failed to enqueue command
throw std::runtime_error("clEnqueueReadBuffer: " + throw std::runtime_error("clEnqueueReadBuffer: " + get_opencl_error(err));
get_opencl_error(err));
} }
mem_out_events_.push_back(std::move(event)); pos += 1;
enqueue_read_buffers(kernel_done, detail::int_list<Is...>{});
} }
template <long I, class T>
void enqueue_read(mem_ref<T>&, std::vector<cl_event>&, size_t&) {
// Nothing to read back if we return references.
}
void enqueue_read_buffers(size_t&, std::vector<cl_event>&,
detail::int_list<>) {
// end of recursion
}
template <long I, long... Is>
void enqueue_read_buffers(size_t& pos, std::vector<cl_event>& events,
detail::int_list<I, Is...>) {
enqueue_read<I>(std::get<I>(results_), events, pos);
enqueue_read_buffers(pos, events, detail::int_list<Is...>{});
}
// handle results if execution result includes a value type
void handle_results() { void handle_results() {
auto actor_facade = auto parent = static_cast<Actor*>(actor_cast<abstract_actor*>(cl_actor_));
static_cast<FacadeType*>(actor_cast<abstract_actor*>(actor_facade_)); auto& map_fun = parent->map_results_;
auto& map_fun = actor_facade->map_results_; auto msg = map_fun ? apply_args(map_fun, detail::get_indices(results_),
auto msg = map_fun ? apply_args(map_fun, results_)
detail::get_indices(result_buffers_), : message_from_results{}(results_);
result_buffers_) promise_.deliver(std::move(msg));
: message_from_results{}(result_buffers_);
get<0>(handle_)->enqueue(actor_facade_, get<1>(handle_), std::move(msg),
nullptr);
} }
std::vector<size_t> lengths_;
response_promise promise_;
strong_actor_ptr cl_actor_;
std::vector<cl_event> mem_in_events_;
std::vector<cl_event> mem_out_events_;
cl_event_ptr callback_;
std::vector<cl_mem_ptr> input_buffers_;
std::vector<cl_mem_ptr> output_buffers_;
std::vector<cl_mem_ptr> scratch_buffers_;
std::tuple<Ts...> results_;
message msg_; // keeps the argument buffers alive for async copy to device
spawn_config config_;
}; };
} // namespace opencl } // namespace opencl
......
...@@ -20,7 +20,7 @@ ...@@ -20,7 +20,7 @@
#ifndef CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP #ifndef CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
#define CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP #define CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
#include "caf/opencl/actor_facade.hpp" #include "caf/opencl/opencl_actor.hpp"
namespace caf { namespace caf {
namespace opencl { namespace opencl {
...@@ -28,13 +28,13 @@ namespace detail { ...@@ -28,13 +28,13 @@ namespace detail {
struct tuple_construct { }; struct tuple_construct { };
template <class... Ts> template <bool PassConfig, class... Ts>
struct cl_spawn_helper { struct cl_spawn_helper {
using impl = opencl::actor_facade<Ts...>; using impl = opencl::opencl_actor<PassConfig, Ts...>;
using map_in_fun = std::function<optional<message> (message&)>; using map_in_fun = typename impl::input_mapping;
using map_out_fun = typename impl::output_mapping; using map_out_fun = typename impl::output_mapping;
actor operator()(actor_config actor_cfg, const opencl::program& p, actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::spawn_config& spawn_cfg, const char* fn, const opencl::spawn_config& spawn_cfg,
Ts&&... xs) const { Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg), return actor_cast<actor>(impl::create(std::move(actor_cfg),
...@@ -42,7 +42,16 @@ struct cl_spawn_helper { ...@@ -42,7 +42,16 @@ struct cl_spawn_helper {
map_in_fun{}, map_out_fun{}, map_in_fun{}, map_out_fun{},
std::move(xs)...)); std::move(xs)...));
} }
actor operator()(actor_config actor_cfg, const opencl::program& p, actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::spawn_config& spawn_cfg,
map_in_fun map_input, Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg),
p, fn, spawn_cfg,
std::move(map_input),
map_out_fun{},
std::move(xs)...));
}
actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::spawn_config& spawn_cfg, const char* fn, const opencl::spawn_config& spawn_cfg,
map_in_fun map_input, map_out_fun map_output, map_in_fun map_input, map_out_fun map_output,
Ts&&... xs) const { Ts&&... xs) const {
...@@ -58,4 +67,4 @@ struct cl_spawn_helper { ...@@ -58,4 +67,4 @@ struct cl_spawn_helper {
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
#endif // CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP #endif // CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
...@@ -22,23 +22,100 @@ ...@@ -22,23 +22,100 @@
#include <vector> #include <vector>
#include "caf/sec.hpp"
#include "caf/opencl/global.hpp" #include "caf/opencl/global.hpp"
#include "caf/opencl/smart_ptr.hpp" #include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
namespace caf { namespace caf {
namespace opencl { namespace opencl {
class program; class program;
class manager; class manager;
template <class T> class mem_ref;
class device;
using device_ptr = intrusive_ptr<device>;
class device { class device : public ref_counted {
public: public:
friend class program; friend class program;
friend class manager; friend class manager;
template <class T> friend class mem_ref;
template <class T, class... Ts>
friend intrusive_ptr<T> caf::make_counted(Ts&&...);
~device();
/// Create an argument for an OpenCL kernel with data placed in global memory.
template <class T>
mem_ref<T> global_argument(const std::vector<T>& data,
cl_mem_flags flags = buffer_type::input_output,
optional<size_t> size = none,
cl_bool blocking = CL_FALSE) {
size_t num_elements = size ? *size : data.size();
size_t buffer_size = sizeof(T) * num_elements;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(), flags,
buffer_size, nullptr);
cl_event_ptr event{v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
queue_.get(), buffer, blocking,
cl_uint{0}, buffer_size, data.data()),
false};
return mem_ref<T>{num_elements, queue_, std::move(buffer), flags,
std::move(event)};
}
/// Create an argument for an OpenCL kernel in global memory without data.
template <class T>
mem_ref<T> scratch_argument(size_t size,
cl_mem_flags flags = buffer_type::scratch_space) {
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(), flags,
sizeof(T) * size, nullptr);
return mem_ref<T>{size, queue_, std::move(buffer), flags, nullptr};
}
template <class T>
expected<mem_ref<T>> copy(mem_ref<T>& mem) {
if (!mem.get())
return make_error(sec::runtime_error, "No memory assigned.");
auto buffer_size = sizeof(T) * mem.size();
cl_event event;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(), mem.access(),
buffer_size, nullptr);
std::vector<cl_event> prev_events;
cl_event e = mem.take_event();
if (e)
prev_events.push_back(e);
auto err = clEnqueueCopyBuffer(queue_.get(), mem.get().get(), buffer,
0, 0, // no offset for now
buffer_size, prev_events.size(),
prev_events.data(), &event);
if (err != CL_SUCCESS)
return make_error(sec::runtime_error, get_opencl_error(err));
// callback to release the previous event
if (e) {
err = clSetEventCallback(event, CL_COMPLETE,
[](cl_event, cl_int, void* data) {
auto tmp = reinterpret_cast<cl_event>(data);
if (tmp)
clReleaseEvent(tmp);
},
e);
if (err != CL_SUCCESS)
return make_error(sec::runtime_error, get_opencl_error(err));
}
// decrements the previous event we used for waiting above
return mem_ref<T>(mem.size(), queue_, std::move(buffer),
mem.access(), {event, false});
}
/// Intialize a new device in a context using a sepcific device_id /// Initialize a new device in a context using a specific device_id
static device create(const context_ptr& context, const device_ptr& device_id, static device_ptr create(const cl_context_ptr& context,
unsigned id); const cl_device_ptr& device_id,
unsigned id);
/// Synchronizes all commands in its queue, waiting for them to finish.
void synchronize();
/// Get the id assigned by caf /// Get the id assigned by caf
inline unsigned get_id() const; inline unsigned get_id() const;
/// Returns device info on CL_DEVICE_ADDRESS_BITS /// Returns device info on CL_DEVICE_ADDRESS_BITS
...@@ -93,21 +170,21 @@ public: ...@@ -93,21 +170,21 @@ public:
inline const std::string& get_name() const; inline const std::string& get_name() const;
private: private:
device(device_ptr device_id, command_queue_ptr queue, context_ptr context, device(cl_device_ptr device_id, cl_command_queue_ptr queue,
unsigned id); cl_context_ptr context, unsigned id);
template <class T> template <class T>
static T info(const device_ptr& device_id, unsigned info_flag) { static T info(const cl_device_ptr& device_id, unsigned info_flag) {
T value; T value;
clGetDeviceInfo(device_id.get(), info_flag, sizeof(T), &value, nullptr); clGetDeviceInfo(device_id.get(), info_flag, sizeof(T), &value, nullptr);
return value; return value;
} }
static std::string info_string(const device_ptr& device_id, static std::string info_string(const cl_device_ptr& device_id,
unsigned info_flag); unsigned info_flag);
device_ptr device_id_; cl_device_ptr device_id_;
command_queue_ptr command_queue_; cl_command_queue_ptr queue_;
context_ptr context_; cl_context_ptr context_;
unsigned id_; unsigned id_;
bool profiling_enabled_; // CL_DEVICE_QUEUE_PROPERTIES bool profiling_enabled_; // CL_DEVICE_QUEUE_PROPERTIES
......
...@@ -50,6 +50,14 @@ enum device_type { ...@@ -50,6 +50,14 @@ enum device_type {
all = CL_DEVICE_TYPE_ALL all = CL_DEVICE_TYPE_ALL
}; };
/// Default values to create OpenCL buffers
enum buffer_type : cl_mem_flags {
input = CL_MEM_READ_WRITE | CL_MEM_HOST_WRITE_ONLY,
input_output = CL_MEM_READ_WRITE,
output = CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY,
scratch_space = CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS
};
std::ostream& operator<<(std::ostream& os, device_type dev); std::ostream& operator<<(std::ostream& os, device_type dev);
device_type device_type_from_ulong(cl_ulong dev); device_type device_type_from_ulong(cl_ulong dev);
...@@ -58,6 +66,9 @@ using dim_vec = detail::limited_vector<size_t, 3>; ...@@ -58,6 +66,9 @@ using dim_vec = detail::limited_vector<size_t, 3>;
std::string get_opencl_error(cl_int err); std::string get_opencl_error(cl_int err);
std::string event_status(cl_event event);
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
......
...@@ -34,7 +34,7 @@ ...@@ -34,7 +34,7 @@
#include "caf/opencl/program.hpp" #include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp" #include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp" #include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/actor_facade.hpp" #include "caf/opencl/opencl_actor.hpp"
#include "caf/opencl/detail/spawn_helper.hpp" #include "caf/opencl/detail/spawn_helper.hpp"
...@@ -45,17 +45,17 @@ class manager : public actor_system::module { ...@@ -45,17 +45,17 @@ class manager : public actor_system::module {
public: public:
friend class program; friend class program;
friend class actor_system; friend class actor_system;
friend command_queue_ptr get_command_queue(uint32_t id); friend cl_command_queue_ptr get_command_queue(uint32_t id);
manager(const manager&) = delete; manager(const manager&) = delete;
manager& operator=(const manager&) = delete; manager& operator=(const manager&) = delete;
/// Get the device with id, which is assigned sequientally. /// Get the device with id, which is assigned sequientally.
const optional<const device&> get_device(size_t dev_id = 0) const; optional<device_ptr> get_device(size_t dev_id = 0) const;
/// Get the first device that satisfies the predicate. /// Get the first device that satisfies the predicate.
/// The predicate should accept a `const device&` and return a bool; /// The predicate should accept a `const device&` and return a bool;
template <class UnaryPredicate> template <class UnaryPredicate>
const optional<const device&> get_device_if(UnaryPredicate p) const { optional<device_ptr> get_device_if(UnaryPredicate p) const {
for (auto& pl : platforms_) { for (auto& pl : platforms_) {
for (auto& dev : pl.get_devices()) { for (auto& dev : pl->get_devices()) {
if (p(dev)) if (p(dev))
return dev; return dev;
} }
...@@ -76,17 +76,32 @@ public: ...@@ -76,17 +76,32 @@ public:
// OpenCL functionality // OpenCL functionality
/// @brief Factory method, that creates a caf::opencl::program
/// reading the source from given @p path.
/// @returns A program object.
program_ptr create_program_from_file(const char* path,
const char* options = nullptr,
uint32_t device_id = 0);
/// @brief Factory method, that creates a caf::opencl::program /// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source. /// from a given @p kernel_source.
/// @returns A program object. /// @returns A program object.
program create_program(const char* kernel_source, program_ptr create_program(const char* kernel_source,
const char* options = nullptr, uint32_t device_id = 0); const char* options = nullptr,
uint32_t device_id = 0);
/// @brief Factory method, that creates a caf::opencl::program
/// reading the source from given @p path.
/// @returns A program object.
program_ptr create_program_from_file(const char* path,
const char* options,
const device_ptr dev);
/// @brief Factory method, that creates a caf::opencl::program /// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source. /// from a given @p kernel_source.
/// @returns A program object. /// @returns A program object.
program create_program(const char* kernel_source, program_ptr create_program(const char* kernel_source,
const char* options, const device& dev); const char* options, const device_ptr dev);
/// Creates a new actor facade for an OpenCL kernel that invokes /// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`. /// the function named `fname` from `prog`.
...@@ -97,12 +112,9 @@ public: ...@@ -97,12 +112,9 @@ public:
opencl::is_opencl_arg<T>::value, opencl::is_opencl_arg<T>::value,
actor actor
>::type >::type
spawn(const opencl::program& prog, spawn(const opencl::program_ptr prog, const char* fname,
const char* fname, const opencl::spawn_config& config, T x, Ts... xs) {
const opencl::spawn_config& config, detail::cl_spawn_helper<false, T, Ts...> f;
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, config, return f(actor_config{system_.dummy_execution_unit()}, prog, fname, config,
std::move(x), std::move(xs)...); std::move(x), std::move(xs)...);
} }
...@@ -117,12 +129,9 @@ public: ...@@ -117,12 +129,9 @@ public:
opencl::is_opencl_arg<T>::value, opencl::is_opencl_arg<T>::value,
actor actor
>::type >::type
spawn(const char* source, spawn(const char* source, const char* fname,
const char* fname, const opencl::spawn_config& config, T x, Ts... xs) {
const opencl::spawn_config& config, detail::cl_spawn_helper<false, T, Ts...> f;
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, config, create_program(source), fname, config,
std::move(x), std::move(xs)...); std::move(x), std::move(xs)...);
...@@ -133,13 +142,11 @@ public: ...@@ -133,13 +142,11 @@ public:
/// @throws std::runtime_error if more than three dimensions are set, /// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed. /// `dims.empty()`, or `clCreateKernel` failed.
template <class Fun, class... Ts> template <class Fun, class... Ts>
actor spawn(const opencl::program& prog, actor spawn(const opencl::program_ptr prog, const char* fname,
const char* fname,
const opencl::spawn_config& config, const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args, std::function<optional<message> (message&)> map_args,
Fun map_result, Fun map_result, Ts... xs) {
Ts... xs) { detail::cl_spawn_helper<false, Ts...> f;
detail::cl_spawn_helper<Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, config, return f(actor_config{system_.dummy_execution_unit()}, prog, fname, config,
std::move(map_args), std::move(map_result), std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...); std::forward<Ts>(xs)...);
...@@ -151,26 +158,145 @@ public: ...@@ -151,26 +158,145 @@ public:
/// <tt>dims.empty()</tt>, a compilation error /// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed. /// occured, or @p clCreateKernel failed.
template <class Fun, class... Ts> template <class Fun, class... Ts>
actor spawn(const char* source, actor spawn(const char* source, const char* fname,
const char* fname,
const opencl::spawn_config& config, const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args, std::function<optional<message> (message&)> map_args,
Fun map_result, Fun map_result, Ts... xs) {
Ts... xs) { detail::cl_spawn_helper<false, Ts...> f;
detail::cl_spawn_helper<Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, config, create_program(source), fname, config,
std::move(map_args), std::move(map_result), std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...); std::forward<Ts>(xs)...);
} }
// --- Only accept the input mapping function ---
/// 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 T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args,
T x, Ts... xs) {
detail::cl_spawn_helper<false, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, config,
std::move(map_args), std::forward<T>(x), 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 T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn(const char* source, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args,
T x, Ts... xs) {
detail::cl_spawn_helper<false, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, config,
std::move(map_args), std::forward<T>(x), std::forward<Ts>(xs)...);
}
// --- Input mapping function accepts config as well ---
/// 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>
typename std::enable_if<
!opencl::is_opencl_arg<Fun>::value,
actor
>::type
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
Fun map_result, Ts... xs) {
detail::cl_spawn_helper<true, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, 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>
typename std::enable_if<
!opencl::is_opencl_arg<Fun>::value,
actor
>::type
spawn(const char* source, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
Fun map_result, Ts... xs) {
detail::cl_spawn_helper<true, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, config,
std::move(map_args), std::move(map_result),
std::forward<Ts>(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 T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
T x, Ts... xs) {
detail::cl_spawn_helper<true, T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, config,
std::move(map_args), std::forward<T>(x), 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 T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn(const char* source, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
T x, Ts... xs) {
detail::cl_spawn_helper<true, T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, config,
std::move(map_args), std::forward<T>(x), std::forward<Ts>(xs)...);
}
protected: protected:
manager(actor_system& sys); manager(actor_system& sys);
~manager() override; ~manager() override;
private: private:
actor_system& system_; actor_system& system_;
std::vector<platform> platforms_; std::vector<platform_ptr> platforms_;
}; };
} // namespace opencl } // namespace opencl
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* 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_MEM_REF_HPP
#define CAF_OPENCL_MEM_REF_HPP
#include <ios>
#include <vector>
#include "caf/sec.hpp"
#include "caf/optional.hpp"
#include "caf/ref_counted.hpp"
#include "caf/opencl/smart_ptr.hpp"
namespace caf {
namespace opencl {
struct msg_adding_event {
msg_adding_event(cl_event_ptr event) : event_(event) {
// nop
}
template <class T, class... Ts>
message operator()(T& x, Ts&... xs) {
return make_message(add_event(std::move(x)), add_event(std::move(xs))...);
}
template <class... Ts>
message operator()(std::tuple<Ts...>& values) {
return apply_args(*this, detail::get_indices(values), values);
}
template <class T>
mem_ref<T> add_event(mem_ref<T> ref) {
ref.set_event(event_);
return std::move(ref);
}
cl_event_ptr event_;
};
class device;
/// A reference type for buffers on a OpenCL devive. Access is not thread safe.
/// Hence, a mem_ref should only be passed to actors sequentially.
template <class T>
class mem_ref {
public:
using value_type = T;
friend struct msg_adding_event;
template <bool PassConfig, class... Ts>
friend class opencl_actor;
friend class device;
expected<std::vector<T>> data(optional<size_t> result_size = none) {
if (!memory_)
return make_error(sec::runtime_error, "No memory assigned.");
if (0 != (access_ & CL_MEM_HOST_NO_ACCESS))
return make_error(sec::runtime_error, "No memory access.");
if (result_size && *result_size > num_elements_)
return make_error(sec::runtime_error, "Buffer has less elements.");
auto num_elements = (result_size ? *result_size : num_elements_);
auto buffer_size = sizeof(T) * num_elements;
std::vector<T> buffer(num_elements);
std::vector<cl_event> prev_events;
if (event_)
prev_events.push_back(event_.get());
cl_event event;
auto err = clEnqueueReadBuffer(queue_.get(), memory_.get(), CL_TRUE,
0, buffer_size, buffer.data(),
static_cast<cl_uint>(prev_events.size()),
prev_events.data(), &event);
if (err != CL_SUCCESS)
return make_error(sec::runtime_error, get_opencl_error(err));
// decrements the previous event we used for waiting above
event_.reset(event, false);
return buffer;
}
void reset() {
num_elements_ = 0;
access_ = CL_MEM_HOST_NO_ACCESS;
memory_.reset();
access_ = 0;
event_.reset();
}
inline const cl_mem_ptr& get() const {
return memory_;
}
inline cl_mem_ptr& get() {
return memory_;
}
inline size_t size() const {
return num_elements_;
}
inline cl_mem_flags access() const {
return access_;
}
mem_ref()
: num_elements_{0},
access_{CL_MEM_HOST_NO_ACCESS},
queue_{nullptr},
event_{nullptr},
memory_{nullptr} {
// nop
}
mem_ref(size_t num_elements, cl_command_queue_ptr queue,
cl_mem_ptr memory, cl_mem_flags access, cl_event_ptr event)
: num_elements_{num_elements},
access_{access},
queue_{queue},
event_{event},
memory_{memory} {
// nop
}
mem_ref(size_t num_elements, cl_command_queue_ptr queue,
cl_mem memory, cl_mem_flags access, cl_event_ptr event)
: num_elements_{num_elements},
access_{access},
queue_{queue},
event_{event},
memory_{memory} {
// nop
}
mem_ref(mem_ref&& other) = default;
mem_ref& operator=(mem_ref<T>&& other) = default;
mem_ref(const mem_ref& other) = default;
mem_ref& operator=(const mem_ref& other) = default;
~mem_ref() {
// nop
}
private:
inline void set_event(cl_event e, bool increment_reference = true) {
event_.reset(e, increment_reference);
}
inline void set_event(cl_event_ptr e) {
event_ = std::move(e);
}
inline cl_event_ptr event() {
return event_;
}
inline cl_event take_event() {
return event_.release();
}
size_t num_elements_;
cl_mem_flags access_;
cl_command_queue_ptr queue_;
cl_event_ptr event_;
cl_mem_ptr memory_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_MEM_REF_HPP
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* 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_OPENCL_ACTOR_HPP
#define CAF_OPENCL_OPENCL_ACTOR_HPP
#include <ostream>
#include <iostream>
#include <algorithm>
#include <stdexcept>
#include "caf/all.hpp"
#include "caf/intrusive_ptr.hpp"
#include "caf/detail/int_list.hpp"
#include "caf/detail/type_list.hpp"
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/command.hpp"
#include "caf/opencl/mem_ref.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 manager;
// signature for the function that is applied to output arguments
template <class List>
struct output_function_sig;
template <class... Ts>
struct output_function_sig<detail::type_list<Ts...>> {
using type = std::function<message (Ts&...)>;
};
// convert to mem_ref
template <class T>
struct to_mem_ref {
using type = mem_ref<T>;
};
template <class T>
struct to_mem_ref<std::vector<T>> {
using type = mem_ref<T>;
};
template <class T>
struct to_mem_ref<T*> {
using type = mem_ref<T>;
};
// derive signature of the command that handles the kernel execution
template <class T, class List>
struct command_sig;
template <class T, class... Ts>
struct command_sig<T, detail::type_list<Ts...>> {
using type = command<T, Ts...>;
};
// derive type for a tuple matching the arguments as mem_refs
template <class List>
struct tuple_type_of;
template <class... Ts>
struct tuple_type_of<detail::type_list<Ts...>> {
using type = std::tuple<Ts...>;
};
template <bool PassConfig, class... Ts>
class opencl_actor : public monitorable_actor {
public:
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_input_type>::type;
using input_mapping = typename std::conditional<PassConfig,
std::function<optional<message> (spawn_config&, message&)>,
std::function<optional<message> (message&)>
>::type;
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_output_type>::type;
using output_mapping = typename output_function_sig<output_types>::type;
using processing_list = typename cl_arg_info_list<arg_types>::type;
using command_type = typename command_sig<opencl_actor, output_types>::type;
typename detail::il_indices<arg_types>::type indices;
using evnt_vec = std::vector<cl_event>;
using mem_vec = std::vector<cl_mem_ptr>;
using len_vec = std::vector<size_t>;
using out_tup = typename tuple_type_of<output_types>::type;
const char* name() const override {
return "OpenCL actor";
}
static actor create(actor_config actor_conf, const program_ptr prog,
const char* kernel_name, const spawn_config& spawn_conf,
input_mapping map_args, output_mapping map_result,
Ts&&... xs) {
if (spawn_conf.dimensions().empty()) {
auto str = "OpenCL kernel needs at least 1 global dimension.";
CAF_LOG_ERROR(str);
throw std::runtime_error(str);
}
auto check_vec = [&](const dim_vec& vec, const char* name) {
if (! vec.empty() && vec.size() != spawn_conf.dimensions().size()) {
std::ostringstream oss;
oss << name << " vector is not empty, but "
<< "its size differs from global dimensions vector's size";
CAF_LOG_ERROR(CAF_ARG(oss.str()));
throw std::runtime_error(oss.str());
}
};
check_vec(spawn_conf.offsets(), "offsets");
check_vec(spawn_conf.local_dimensions(), "local dimensions");
auto& sys = actor_conf.host->system();
auto itr = prog->available_kernels_.find(kernel_name);
if (itr == prog->available_kernels_.end()) {
cl_kernel_ptr kernel;
kernel.reset(v2get(CAF_CLF(clCreateKernel), prog->program_.get(),
kernel_name),
false);
return make_actor<opencl_actor, actor>(sys.next_actor_id(), sys.node(),
&sys, std::move(actor_conf),
prog, kernel, spawn_conf,
std::move(map_args),
std::move(map_result),
std::forward_as_tuple(xs...));
}
return make_actor<opencl_actor, actor>(sys.next_actor_id(), sys.node(),
&sys, std::move(actor_conf),
prog, itr->second, spawn_conf,
std::move(map_args),
std::move(map_result),
std::forward_as_tuple(xs...));
}
void enqueue(strong_actor_ptr sender, message_id mid,
message content, response_promise promise) {
CAF_PUSH_AID(id());
CAF_LOG_TRACE("");
if (!map_arguments(content))
return;
if (!content.match_elements(input_types{})) {
CAF_LOG_ERROR("Message types do not match the expected signature.");
return;
}
auto hdl = std::make_tuple(sender, mid.response_id());
evnt_vec events;
mem_vec input_buffers;
mem_vec output_buffers;
mem_vec scratch_buffers;
len_vec result_lengths;
out_tup result;
add_kernel_arguments(events, // accumulate events for execution
input_buffers, // opencl buffers included in in msg
output_buffers, // opencl buffers included in out msg
scratch_buffers, // opencl only used here
result, // tuple to save the output values
result_lengths, // size of buffers to read back
content, // message content
indices); // enable extraction of types from msg
auto cmd = make_counted<command_type>(
std::move(promise),
actor_cast<strong_actor_ptr>(this),
std::move(events),
std::move(input_buffers),
std::move(output_buffers),
std::move(scratch_buffers),
std::move(result_lengths),
std::move(content),
std::move(result),
config_
);
cmd->enqueue();
}
void enqueue(mailbox_element_ptr ptr, execution_unit* eu) override {
CAF_ASSERT(ptr != nullptr);
CAF_LOG_TRACE(CAF_ARG(*ptr));
response_promise promise{eu, ctrl(), *ptr};
enqueue(ptr->sender, ptr->mid, ptr->move_content_to_message(),
std::move(promise));
}
void enqueue(strong_actor_ptr sender, message_id mid,
message content, execution_unit* host) override {
CAF_LOG_TRACE("");
enqueue(make_mailbox_element(sender, mid, {}, std::move(content)), host);
}
opencl_actor(actor_config actor_conf, const program_ptr prog,
cl_kernel_ptr kernel, spawn_config spawn_conf,
input_mapping map_args, output_mapping map_result,
std::tuple<Ts...> xs)
: monitorable_actor(actor_conf),
kernel_(std::move(kernel)),
program_(prog->program_),
context_(prog->context_),
queue_(prog->queue_),
config_(std::move(spawn_conf)),
map_args_(std::move(map_args)),
map_results_(std::move(map_result)),
kernel_signature_(std::move(xs)) {
CAF_LOG_TRACE(CAF_ARG(this->id()));
default_length_ = std::accumulate(std::begin(config_.dimensions()),
std::end(config_.dimensions()),
size_t{1},
std::multiplies<size_t>{});
}
void add_kernel_arguments(evnt_vec&, mem_vec&, mem_vec&, mem_vec&,
out_tup&, len_vec&, message&, detail::int_list<>) {
// nop
}
/// The separation into input, output and scratch buffers is required to
/// access the related memory handles later on. The scratch and input handles
/// are saved to prevent deletion before the kernel finished execution.
template <long I, long... Is>
void add_kernel_arguments(evnt_vec& events, mem_vec& inputs, mem_vec& outputs,
mem_vec& scratch, out_tup& result, len_vec& lengths,
message& msg, detail::int_list<I, Is...>) {
using arg_type = typename caf::detail::tl_at<processing_list,I>::type;
create_buffer<I, arg_type::in_pos, arg_type::out_pos>(
std::get<I>(kernel_signature_), events, lengths, inputs,
outputs, scratch, result, msg
);
add_kernel_arguments(events, inputs, outputs, scratch, result, lengths, msg,
detail::int_list<Is...>{});
}
// Two functions to handle `in` arguments: val and mref
template <long I, int InPos, int OutPos, class T>
void create_buffer(const in<T, val>&, evnt_vec& events, len_vec&,
mem_vec& inputs, mem_vec&, mem_vec&, out_tup&,
message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
using container_type = std::vector<value_type>;
auto& container = msg.get_as<container_type>(InPos);
auto len = container.size();
size_t num_bytes = sizeof(value_type) * len;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
size_t{CL_MEM_READ_WRITE}, num_bytes, nullptr);
auto event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
queue_.get(), buffer, 0u, // --> CL_FALSE,
0u, num_bytes, container.data());
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&buffer));
events.push_back(event);
inputs.emplace_back(buffer, false);
}
template <long I, int InPos, int OutPos, class T>
void create_buffer(const in<T, mref>&, evnt_vec& events, len_vec&, mem_vec&,
mem_vec&, mem_vec&, out_tup&, message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
using container_type = mem_ref<value_type>;
auto container = msg.get_as<container_type>(InPos);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&container.get()));
auto event = container.take_event();
if (event)
events.push_back(event);
}
// Four functions to handle `in_out` arguments:
// val->val, val->mref, mref->val, mref->mref
template <long I, int InPos, int OutPos, class T>
void create_buffer(const in_out<T,val,val>&, evnt_vec& events,
len_vec& lengths, mem_vec&, mem_vec& outputs,
mem_vec&, out_tup&, message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
using container_type = std::vector<value_type>;
auto& container = msg.get_as<container_type>(InPos);
auto len = container.size();
size_t num_bytes = sizeof(value_type) * len;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
size_t{CL_MEM_READ_WRITE}, num_bytes, nullptr);
auto event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
queue_.get(), buffer, 0u, // --> CL_FALSE,
0u, num_bytes, container.data());
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&buffer));
lengths.push_back(len);
events.push_back(event);
outputs.emplace_back(buffer, false);
}
template <long I, int InPos, int OutPos, class T>
void create_buffer(const in_out<T,val,mref>&, evnt_vec& events, len_vec&,
mem_vec&, mem_vec&, mem_vec&, out_tup& result,
message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
using container_type = std::vector<value_type>;
auto& container = msg.get_as<container_type>(InPos);
auto len = container.size();
size_t num_bytes = sizeof(value_type) * len;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
size_t{CL_MEM_READ_WRITE}, num_bytes, nullptr);
auto event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
queue_.get(), buffer, 0u, // --> CL_FALSE,
0u, num_bytes, container.data());
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&buffer));
events.push_back(event);
std::get<OutPos>(result) = mem_ref<value_type>{
len, queue_, cl_mem_ptr{buffer, false},
size_t{CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY}, nullptr
};
}
template <long I, int InPos, int OutPos, class T>
void create_buffer(const in_out<T,mref,val>&, evnt_vec& events,
len_vec& lengths, mem_vec&, mem_vec& outputs,
mem_vec&, out_tup&, message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
using container_type = mem_ref<value_type>;
auto container = msg.get_as<container_type>(InPos);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&container.get()));
auto event = container.take_event();
if (event)
events.push_back(event);
lengths.push_back(container.size());
outputs.push_back(container.get());
}
template <long I, int InPos, int OutPos, class T>
void create_buffer(const in_out<T,mref,mref>&, evnt_vec& events, len_vec&,
mem_vec&, mem_vec&, mem_vec&, out_tup& result,
message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
using container_type = mem_ref<value_type>;
auto container = msg.get_as<container_type>(InPos);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&container.get()));
auto event = container.take_event();
if (event)
events.push_back(event);
std::get<OutPos>(result) = container;
}
// Two functions to handle `out` arguments: val and mref
template <long I, int InPos, int OutPos, class T>
void create_buffer(const out<T,val>& wrapper, evnt_vec&, len_vec& lengths,
mem_vec&, mem_vec& outputs, mem_vec&, out_tup&,
message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
auto len = argument_length(wrapper, msg, default_length_);
auto num_bytes = sizeof(value_type) * len;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
size_t{CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY},
num_bytes, nullptr);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&buffer));
outputs.emplace_back(buffer, false);
lengths.push_back(len);
}
template <long I, int InPos, int OutPos, class T>
void create_buffer(const out<T,mref>& wrapper, evnt_vec&, len_vec&,
mem_vec&, mem_vec&, mem_vec&, out_tup& result,
message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
auto len = argument_length(wrapper, msg, default_length_);
auto num_bytes = sizeof(value_type) * len;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
size_t{CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY},
num_bytes, nullptr);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&buffer));
std::get<OutPos>(result) = mem_ref<value_type>{
len, queue_, {buffer, false},
size_t{CL_MEM_READ_WRITE | CL_MEM_HOST_READ_ONLY}, nullptr
};
}
// One function to handle `scratch` buffers
template <long I, int InPos, int OutPos, class T>
void create_buffer(const scratch<T>& wrapper, evnt_vec&, len_vec&,
mem_vec&, mem_vec&, mem_vec& scratch,
out_tup&, message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
auto len = argument_length(wrapper, msg, default_length_);
auto num_bytes = sizeof(value_type) * len;
auto buffer = v2get(CAF_CLF(clCreateBuffer), context_.get(),
size_t{CL_MEM_READ_WRITE | CL_MEM_HOST_NO_ACCESS},
num_bytes, nullptr);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
sizeof(cl_mem), static_cast<const void*>(&buffer));
scratch.emplace_back(buffer, false);
}
// One functions to handle `local` arguments
template <long I, int InPos, int OutPos, class T>
void create_buffer(const local<T>& wrapper, evnt_vec&, len_vec&,
mem_vec&, mem_vec&, mem_vec&, out_tup&,
message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
auto len = wrapper(msg);
auto num_bytes = sizeof(value_type) * len;
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
num_bytes, nullptr);
}
// Two functions to handle `priv` arguments: val and hidden
template <long I, int InPos, int OutPos, class T>
void create_buffer(const priv<T, val>&, evnt_vec&, len_vec&,
mem_vec&, mem_vec&, mem_vec&, out_tup&, message& msg) {
using value_type = typename detail::tl_at<unpacked_types, I>::type;
auto value_size = sizeof(value_type);
auto& value = msg.get_as<value_type>(InPos);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
value_size, static_cast<const void*>(&value));
}
template <long I, int InPos, int OutPos, class T>
void create_buffer(const priv<T, hidden>& wrapper, evnt_vec&, len_vec&,
mem_vec&, mem_vec&, mem_vec&, out_tup&, message& msg) {
auto value_size = sizeof(T);
auto value = wrapper(msg);
v1callcl(CAF_CLF(clSetKernelArg), kernel_.get(), static_cast<unsigned>(I),
value_size, static_cast<const void*>(&value));
}
/// Helper function to calculate the elements in a buffer from in and out
/// argument wrappers.
template <class Fun>
size_t argument_length(Fun& f, message& m, size_t fallback) {
auto length = f(m);
return length && (*length > 0) ? *length : fallback;
}
// Map function requires only the message as argument
template <bool Q = PassConfig>
typename std::enable_if<!Q, bool>::type map_arguments(message& content) {
if (map_args_) {
auto mapped = map_args_(content);
if (!mapped) {
CAF_LOG_ERROR("Mapping argumentes failed.");
return false;
}
content = std::move(*mapped);
}
return true;
}
// Map function requires reference to config as well as the message
template <bool Q = PassConfig>
typename std::enable_if<Q, bool>::type map_arguments(message& content) {
if (map_args_) {
auto mapped = map_args_(config_, content);
if (!mapped) {
CAF_LOG_ERROR("Mapping argumentes failed.");
return false;
}
content = std::move(*mapped);
}
return true;
}
cl_kernel_ptr kernel_;
cl_program_ptr program_;
cl_context_ptr context_;
cl_command_queue_ptr queue_;
spawn_config config_;
input_mapping map_args_;
output_mapping map_results_;
std::tuple<Ts...> kernel_signature_;
size_t default_length_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_OPENCL_ACTOR_HPP
...@@ -20,40 +20,50 @@ ...@@ -20,40 +20,50 @@
#ifndef CAF_OPENCL_PLATFORM_HPP #ifndef CAF_OPENCL_PLATFORM_HPP
#define CAF_OPENCL_PLATFORM_HPP #define CAF_OPENCL_PLATFORM_HPP
#include <caf/opencl/device.hpp> #include "caf/ref_counted.hpp"
#include "caf/opencl/device.hpp"
namespace caf { namespace caf {
namespace opencl { namespace opencl {
class platform { class platform;
using platform_ptr = intrusive_ptr<platform>;
class platform : public ref_counted {
public: public:
friend class program; friend class program;
template <class T, class... Ts>
friend intrusive_ptr<T> caf::make_counted(Ts&&...);
inline const std::vector<device>& get_devices() const; inline const std::vector<device_ptr>& get_devices() const;
inline const std::string& get_name() const; inline const std::string& get_name() const;
inline const std::string& get_vendor() const; inline const std::string& get_vendor() const;
inline const std::string& get_version() const; inline const std::string& get_version() const;
static platform create(cl_platform_id platform_id, unsigned start_id); static platform_ptr create(cl_platform_id platform_id, unsigned start_id);
private: private:
platform(cl_platform_id platform_id, context_ptr context, platform(cl_platform_id platform_id, cl_context_ptr context,
std::string name, std::string vendor, std::string version, std::string name, std::string vendor, std::string version,
std::vector<device> devices); std::vector<device_ptr> devices);
~platform();
static std::string platform_info(cl_platform_id platform_id, static std::string platform_info(cl_platform_id platform_id,
unsigned info_flag); unsigned info_flag);
cl_platform_id platform_id_; cl_platform_id platform_id_;
context_ptr context_; cl_context_ptr context_;
std::string name_; std::string name_;
std::string vendor_; std::string vendor_;
std::string version_; std::string version_;
std::vector<device> devices_; std::vector<device_ptr> devices_;
}; };
/******************************************************************************\ /******************************************************************************\
* implementation of inline member functions * * implementation of inline member functions *
\******************************************************************************/ \******************************************************************************/
inline const std::vector<device>& platform::get_devices() const { inline const std::vector<device_ptr>& platform::get_devices() const {
return devices_; return devices_;
} }
...@@ -73,4 +83,4 @@ inline const std::string& platform::get_version() const { ...@@ -73,4 +83,4 @@ inline const std::string& platform::get_version() const {
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
#endif // CAF_OPENCL_PLTFORM_HPP #endif // CAF_OPENCL_PLATFORM_HPP
...@@ -23,6 +23,8 @@ ...@@ -23,6 +23,8 @@
#include <map> #include <map>
#include <memory> #include <memory>
#include "caf/ref_counted.hpp"
#include "caf/opencl/device.hpp" #include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp" #include "caf/opencl/global.hpp"
#include "caf/opencl/smart_ptr.hpp" #include "caf/opencl/smart_ptr.hpp"
...@@ -33,21 +35,31 @@ namespace opencl { ...@@ -33,21 +35,31 @@ namespace opencl {
template <class... Ts> template <class... Ts>
class actor_facade; class actor_facade;
class program;
using program_ptr = intrusive_ptr<program>;
/// @brief A wrapper for OpenCL's cl_program. /// @brief A wrapper for OpenCL's cl_program.
class program { class program : public ref_counted {
public: public:
friend class manager; friend class manager;
template <class... Ts> template <class... Ts>
friend class actor_facade; friend class actor_facade;
template <bool PassConfig, class... Ts>
friend class opencl_actor;
template <class T, class... Ts>
friend intrusive_ptr<T> caf::make_counted(Ts&&...);
private: private:
program(context_ptr context, command_queue_ptr queue, program_ptr prog, program(cl_context_ptr context, cl_command_queue_ptr queue,
std::map<std::string,kernel_ptr> available_kernels); cl_program_ptr prog,
std::map<std::string,cl_kernel_ptr> available_kernels);
~program();
context_ptr context_; cl_context_ptr context_;
program_ptr program_; cl_program_ptr program_;
command_queue_ptr queue_; cl_command_queue_ptr queue_;
std::map<std::string,kernel_ptr> available_kernels_; std::map<std::string,cl_kernel_ptr> available_kernels_;
}; };
} // namespace opencl } // namespace opencl
......
...@@ -37,20 +37,20 @@ ...@@ -37,20 +37,20 @@
} /* namespace opencl */ \ } /* namespace opencl */ \
} // namespace caf } // namespace caf
CAF_OPENCL_PTR_ALIAS(mem_ptr, cl_mem, clRetainMemObject, clReleaseMemObject) CAF_OPENCL_PTR_ALIAS(cl_mem_ptr, cl_mem, clRetainMemObject, clReleaseMemObject)
CAF_OPENCL_PTR_ALIAS(event_ptr, cl_event, clRetainEvent, clReleaseEvent) CAF_OPENCL_PTR_ALIAS(cl_event_ptr, cl_event, clRetainEvent, clReleaseEvent)
CAF_OPENCL_PTR_ALIAS(kernel_ptr, cl_kernel, clRetainKernel, clReleaseKernel) CAF_OPENCL_PTR_ALIAS(cl_kernel_ptr, cl_kernel, clRetainKernel, clReleaseKernel)
CAF_OPENCL_PTR_ALIAS(context_ptr, cl_context, clRetainContext, clReleaseContext) CAF_OPENCL_PTR_ALIAS(cl_context_ptr, cl_context, clRetainContext, clReleaseContext)
CAF_OPENCL_PTR_ALIAS(program_ptr, cl_program, clRetainProgram, clReleaseProgram) CAF_OPENCL_PTR_ALIAS(cl_program_ptr, cl_program, clRetainProgram, clReleaseProgram)
CAF_OPENCL_PTR_ALIAS(device_ptr, cl_device_id, CAF_OPENCL_PTR_ALIAS(cl_device_ptr, cl_device_id,
clRetainDeviceDummy, clReleaseDeviceDummy) clRetainDeviceDummy, clReleaseDeviceDummy)
CAF_OPENCL_PTR_ALIAS(command_queue_ptr, cl_command_queue, CAF_OPENCL_PTR_ALIAS(cl_command_queue_ptr, cl_command_queue,
clRetainCommandQueue, clReleaseCommandQueue) clRetainCommandQueue, clReleaseCommandQueue)
#endif // CAF_OPENCL_SMART_PTR_HPP #endif // CAF_OPENCL_SMART_PTR_HPP
...@@ -45,6 +45,12 @@ public: ...@@ -45,6 +45,12 @@ public:
// nop // nop
} }
spawn_config(const spawn_config&) = default;
spawn_config(spawn_config&&) = default;
spawn_config& operator=(const spawn_config&) = default;
spawn_config& operator=(spawn_config&&) = default;
const opencl::dim_vec& dimensions() const { const opencl::dim_vec& dimensions() const {
return dims_; return dims_;
} }
...@@ -58,10 +64,9 @@ public: ...@@ -58,10 +64,9 @@ public:
} }
private: private:
const opencl::dim_vec dims_; opencl::dim_vec dims_;
const opencl::dim_vec offset_; opencl::dim_vec offset_;
const opencl::dim_vec local_dims_; opencl::dim_vec local_dims_;
}; };
} // namespace opencl } // namespace opencl
......
cmake_minimum_required(VERSION 2.8) cmake_minimum_required(VERSION 2.8)
project(caf_examples_opencl CXX) project(caf_examples_opencl CXX)
if(OpenCL_LIBRARIES) if(OpenCL_LIBRARIES AND NOT CAF_NO_EXAMPLES)
add_custom_target(opencl_examples) add_custom_target(opencl_examples)
include_directories(${LIBCAF_INCLUDE_DIRS}) include_directories(${LIBCAF_INCLUDE_DIRS})
if(${CMAKE_SYSTEM_NAME} MATCHES "Window") if(${CMAKE_SYSTEM_NAME} MATCHES "Window")
...@@ -21,4 +21,5 @@ if(OpenCL_LIBRARIES) ...@@ -21,4 +21,5 @@ if(OpenCL_LIBRARIES)
endmacro() endmacro()
add(proper_matrix .) add(proper_matrix .)
add(simple_matrix .) add(simple_matrix .)
add(scan .)
endif() endif()
...@@ -44,17 +44,16 @@ constexpr const char* kernel_name = "matrix_mult"; ...@@ -44,17 +44,16 @@ constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2 // opencl kernel, multiplies matrix1 and matrix2
// last parameter is, by convention, the output parameter // last parameter is, by convention, the output parameter
constexpr const char* kernel_source = R"__( constexpr const char* kernel_source = R"__(
__kernel void matrix_mult(__global float* matrix1, kernel void matrix_mult(global const float* matrix1,
__global float* matrix2, global const float* matrix2,
__global float* output) { global float* output) {
// we only use square matrices, hence: width == height // we only use square matrices, hence: width == height
size_t size = get_global_size(0); // == get_global_size_(1); size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0); size_t x = get_global_id(0);
size_t y = get_global_id(1); size_t y = get_global_id(1);
float result = 0; float result = 0;
for (size_t idx = 0; idx < size; ++idx) { for (size_t idx = 0; idx < size; ++idx)
result += matrix1[idx + y * size] * matrix2[x + idx * size]; result += matrix1[idx + y * size] * matrix2[x + idx * size];
}
output[x+y*size] = result; output[x+y*size] = result;
} }
)__"; )__";
...@@ -83,7 +82,6 @@ public: ...@@ -83,7 +82,6 @@ public:
explicit square_matrix(fvec d) : data_(move(d)) { explicit square_matrix(fvec d) : data_(move(d)) {
assert(data_.size() == num_elements); assert(data_.size() == num_elements);
} }
inline float& operator()(size_t column, size_t row) { inline float& operator()(size_t column, size_t row) {
...@@ -127,13 +125,13 @@ string to_string(const square_matrix<Size>& m) { ...@@ -127,13 +125,13 @@ string to_string(const square_matrix<Size>& m) {
// to annouce the square_matrix to libcaf, we have to define these operators // to annouce the square_matrix to libcaf, we have to define these operators
template<size_t Size> template<size_t Size>
inline bool operator==(const square_matrix<Size>& lhs, inline bool operator==(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) { const square_matrix<Size>& rhs) {
return equal(lhs.begin(), lhs.end(), rhs.begin()); return equal(lhs.begin(), lhs.end(), rhs.begin());
} }
template<size_t Size> template<size_t Size>
inline bool operator!=(const square_matrix<Size>& lhs, inline bool operator!=(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) { const square_matrix<Size>& rhs) {
return !(lhs == rhs); return !(lhs == rhs);
} }
...@@ -149,14 +147,12 @@ void multiplier(event_based_actor* self) { ...@@ -149,14 +147,12 @@ void multiplier(event_based_actor* self) {
// print "source" matrix // print "source" matrix
cout << "calculating square of matrix:" << endl cout << "calculating square of matrix:" << endl
<< to_string(m1) << endl; << to_string(m1) << endl;
auto unbox_args = [](message& msg) -> optional<message> { auto unbox_args = [](message& msg) -> optional<message> {
return msg.apply( return msg.apply([](matrix_type& lhs, matrix_type& rhs) {
[](matrix_type& lhs, matrix_type& rhs) { return make_message(std::move(lhs.data()), std::move(rhs.data()));
return make_message(std::move(lhs.data()), std::move(rhs.data())); });
}
);
}; };
auto box_res = [] (fvec& result) -> message { auto box_res = [] (fvec& result) -> message {
...@@ -176,11 +172,13 @@ void multiplier(event_based_actor* self) { ...@@ -176,11 +172,13 @@ void multiplier(event_based_actor* self) {
// 5th arg: converts the ouptut vector back to a matrix that is then // 5th arg: converts the ouptut vector back to a matrix that is then
// used as response message // used as response message
// from 6 : a description of the kernel signature using in/out/in_out classes // from 6 : a description of the kernel signature using in/out/in_out classes
// with the argument type packed in vectors // with the argument type. Since the actor always expects input
// arguments for global memory to be contained in vectors,
// the vector is omitted here.
auto worker = mngr.spawn(kernel_source, kernel_name, auto worker = mngr.spawn(kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}}, spawn_config{dim_vec{matrix_size, matrix_size}},
unbox_args, box_res, unbox_args, box_res,
in<fvec>{}, in<fvec>{}, out<fvec>{}); in<float>{}, in<float>{}, out<float>{});
// send both matrices to the actor and // send both matrices to the actor and
// wait for results in form of a matrix_type // wait for results in form of a matrix_type
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* 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. *
******************************************************************************/
#include <vector>
#include <random>
#include <iomanip>
#include <numeric>
#include <iostream>
#include "caf/all.hpp"
#include "caf/opencl/all.hpp"
using namespace std;
using namespace caf;
using namespace caf::opencl;
using caf::detail::limited_vector;
namespace {
using uval = unsigned;
using uvec = std::vector<uval>;
using uref = mem_ref<uval>;
constexpr const size_t problem_size = 23;
constexpr const char* kernel_name_1 = "phase_1";
constexpr const char* kernel_name_2 = "phase_2";
constexpr const char* kernel_name_3 = "phase_3";
// opencl kernel, exclusive scan
// last parameter is, by convention, the output parameter
constexpr const char* kernel_source = R"__(
/// Global exclusive scan, phase 1. From:
/// - http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html
kernel void phase_1(global uint* restrict data,
global uint* restrict increments,
local uint* tmp, uint len) {
const uint thread = get_local_id(0);
const uint block = get_group_id(0);
const uint threads_per_block = get_local_size(0);
const uint elements_per_block = threads_per_block * 2;
const uint global_offset = block * elements_per_block;
const uint n = elements_per_block;
uint offset = 1;
// A (2 lines) --> load input into shared memory
tmp[2 * thread] = (global_offset + (2 * thread) < len)
? data[global_offset + (2 * thread)] : 0;
tmp[2 * thread + 1] = (global_offset + (2 * thread + 1) < len)
? data[global_offset + (2 * thread + 1)] : 0;
// build sum in place up the tree
for (uint d = n >> 1; d > 0; d >>= 1) {
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// B (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
tmp[bi] += tmp[ai];
}
offset *= 2;
}
// C (2 lines) --> clear the last element
if (thread == 0) {
increments[block] = tmp[n - 1];
tmp[n - 1] = 0;
}
// traverse down tree & build scan
for (uint d = 1; d < n; d *= 2) {
offset >>= 1;
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// D (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
uint t = tmp[ai];
tmp[ai] = tmp[bi];
tmp[bi] += t;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
// E (2 line) --> write results to device memory
if (global_offset + (2 * thread) < len)
data[global_offset + (2 * thread)] = tmp[2 * thread];
if (global_offset + (2 * thread + 1) < len)
data[global_offset + (2 * thread + 1)] = tmp[2 * thread + 1];
}
/// Global exclusive scan, phase 2.
kernel void phase_2(global uint* restrict data, // not used ...
global uint* restrict increments,
uint len) {
local uint tmp[2048];
uint thread = get_local_id(0);
uint offset = 1;
const uint n = 2048;
// A (2 lines) --> load input into shared memory
tmp[2 * thread] = (2 * thread < len) ? increments[2 * thread] : 0;
tmp[2 * thread + 1] = (2 * thread + 1 < len) ? increments[2 * thread + 1] : 0;
// build sum in place up the tree
for (uint d = n >> 1; d > 0; d >>= 1) {
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// B (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
tmp[bi] += tmp[ai];
}
offset *= 2;
}
// C (2 lines) --> clear the last element
if (thread == 0)
tmp[n - 1] = 0;
// traverse down tree & build scan
for (uint d = 1; d < n; d *= 2) {
offset >>= 1;
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// D (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
uint t = tmp[ai];
tmp[ai] = tmp[bi];
tmp[bi] += t;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
// E (2 line) --> write results to device memory
if (2 * thread < len) increments[2 * thread] = tmp[2 * thread];
if (2 * thread + 1 < len) increments[2 * thread + 1] = tmp[2 * thread + 1];
}
kernel void phase_3(global uint* restrict data,
global uint* restrict increments,
uint len) {
const uint thread = get_local_id(0);
const uint block = get_group_id(0);
const uint threads_per_block = get_local_size(0);
const uint elements_per_block = threads_per_block * 2;
const uint global_offset = block * elements_per_block;
// add the appropriate value to each block
uint ai = 2 * thread;
uint bi = 2 * thread + 1;
uint ai_global = ai + global_offset;
uint bi_global = bi + global_offset;
uint increment = increments[block];
if (ai_global < len) data[ai_global] += increment;
if (bi_global < len) data[bi_global] += increment;
}
)__";
} // namespace <anonymous>
// Allow sending of unserializable references
CAF_ALLOW_UNSAFE_MESSAGE_TYPE(uref);
template <class T, class E = typename enable_if<is_integral<T>::value>::type>
T round_up(T numToRound, T multiple) {
return ((numToRound + multiple - 1) / multiple) * multiple;
}
int main() {
actor_system_config cfg;
cfg.load<opencl::manager>()
.add_message_type<uvec>("uint_vector");
actor_system system{cfg};
cout << "Calculating exclusive scan of '" << problem_size
<< "' values." << endl;
// ---- create data ----
uvec values;
values.reserve(problem_size);
random_device rd; //Will be used to obtain a seed for the random number engine
mt19937 gen(rd()); //Standard mersenne_twister_engine seeded with rd()
uniform_int_distribution<uval> val_gen(0, 1023);
for (size_t i = 0; i < problem_size; ++i)
values.emplace_back(val_gen(gen));
// ---- find device ----
auto& mngr = system.opencl_manager();
//
string prefix = "GeForce";
auto opt = mngr.get_device_if([&](const device_ptr dev) {
auto& name = dev->get_name();
return equal(begin(prefix), end(prefix), begin(name));
});
if (!opt) {
cout << "No device starting with '" << prefix << "' found. "
<< "Will try the first OpenCL device available." << endl;
opt = mngr.get_device();
}
if (!opt) {
cerr << "Not OpenCL device available." << endl;
return 0;
} else {
cerr << "Found device '" << (*opt)->get_name() << "'." << endl;
}
{
// ---- general ----
auto dev = move(*opt);
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{system};
// ---- config parameters ----
auto half_block = dev->get_max_work_group_size() / 2;
auto get_size = [half_block](size_t n) -> size_t {
return round_up((n + 1) / 2, half_block);
};
auto nd_conf = [half_block, get_size](size_t dim) {
return spawn_config{dim_vec{get_size(dim)}, {}, dim_vec{half_block}};
};
auto reduced_ref = [&](const uref&, uval n) {
// calculate number of groups from the group size from the values size
return size_t{get_size(n) / half_block};
};
// default nd-range
auto ndr = spawn_config{dim_vec{half_block}, {}, dim_vec{half_block}};
// ---- scan actors ----
auto phase1 = mngr.spawn(
prog, kernel_name_1, ndr,
[nd_conf](spawn_config& conf, message& msg) -> optional<message> {
return msg.apply([&](uvec& vec) {
auto size = vec.size();
conf = nd_conf(size);
return make_message(std::move(vec), static_cast<uval>(size));
});
},
in_out<uval, val, mref>{},
out<uval,mref>{reduced_ref},
local<uval>{half_block * 2},
priv<uval, val>{}
);
auto phase2 = mngr.spawn(
prog, kernel_name_2, ndr,
[nd_conf](spawn_config& conf, message& msg) -> optional<message> {
return msg.apply([&](uref& data, uref& incs) {
auto size = incs.size();
conf = nd_conf(size);
return make_message(move(data), move(incs), static_cast<uval>(size));
});
},
in_out<uval,mref,mref>{},
in_out<uval,mref,mref>{},
priv<uval, val>{}
);
auto phase3 = mngr.spawn(
prog, kernel_name_3, ndr,
[nd_conf](spawn_config& conf, message& msg) -> optional<message> {
return msg.apply([&](uref& data, uref& incs) {
auto size = incs.size();
conf = nd_conf(size);
return make_message(move(data), move(incs), static_cast<uval>(size));
});
},
in_out<uval,mref,val>{},
in<uval,mref>{},
priv<uval, val>{}
);
// ---- composed scan actor ----
auto scanner = phase3 * phase2 * phase1;
// ---- scan the data ----
self->send(scanner, values);
self->receive(
[&](const uvec& results) {
cout << "Received results." << endl;
cout << " index | values | scan " << endl
<< "-------+--------+--------" << endl;
for (size_t i = 0; i < problem_size; ++i)
cout << setw(6) << i << " | " << setw(6) << values[i] << " | "
<< setw(6) << results[i] << endl;
}
);
}
system.await_all_actors_done();
return 0;
}
...@@ -41,17 +41,16 @@ constexpr const char* kernel_name = "matrix_mult"; ...@@ -41,17 +41,16 @@ constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2 // opencl kernel, multiplies matrix1 and matrix2
// last parameter is, by convention, the output parameter // last parameter is, by convention, the output parameter
constexpr const char* kernel_source = R"__( constexpr const char* kernel_source = R"__(
__kernel void matrix_mult(__global float* matrix1, kernel void matrix_mult(global const float* matrix1,
__global float* matrix2, global const float* matrix2,
__global float* output) { global float* output) {
// we only use square matrices, hence: width == height // we only use square matrices, hence: width == height
size_t size = get_global_size(0); // == get_global_size_(1); size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0); size_t x = get_global_id(0);
size_t y = get_global_id(1); size_t y = get_global_id(1);
float result = 0; float result = 0;
for (size_t idx = 0; idx < size; ++idx) { for (size_t idx = 0; idx < size; ++idx)
result += matrix1[idx + y * size] * matrix2[x + idx * size]; result += matrix1[idx + y * size] * matrix2[x + idx * size];
}
output[x+y*size] = result; output[x+y*size] = result;
} }
)__"; )__";
...@@ -62,7 +61,7 @@ void print_as_matrix(const fvec& matrix) { ...@@ -62,7 +61,7 @@ void print_as_matrix(const fvec& matrix) {
for (size_t column = 0; column < matrix_size; ++column) { for (size_t column = 0; column < matrix_size; ++column) {
for (size_t row = 0; row < matrix_size; ++row) { for (size_t row = 0; row < matrix_size; ++row) {
cout << fixed << setprecision(2) << setw(9) cout << fixed << setprecision(2) << setw(9)
<< matrix[row + column * matrix_size]; << matrix[row + column * matrix_size];
} }
cout << endl; cout << endl;
} }
...@@ -92,11 +91,13 @@ void multiplier(event_based_actor* self) { ...@@ -92,11 +91,13 @@ void multiplier(event_based_actor* self) {
// - offsets for global dimensions (optional) // - offsets for global dimensions (optional)
// - local dimensions (optional) // - local dimensions (optional)
// 4th to Nth arg: the kernel signature described by in/out/in_out classes // 4th to Nth arg: the kernel signature described by in/out/in_out classes
// that contain the argument type in their template, requires vectors // that contain the argument type in their template. Since the actor
// expects its arguments for global memory to be passed in vectors,
// the vector type is omitted for brevity.
auto worker = self->system().opencl_manager().spawn( auto worker = self->system().opencl_manager().spawn(
kernel_source, kernel_name, kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}}, spawn_config{dim_vec{matrix_size, matrix_size}},
in<fvec>{}, in<fvec>{}, out<fvec>{} in<float>{}, in<float>{}, out<float>{}
); );
// send both matrices to the actor and wait for a result // send both matrices to the actor and wait for a result
self->request(worker, chrono::seconds(5), move(m1), move(m2)).then( self->request(worker, chrono::seconds(5), move(m1), move(m2)).then(
......
...@@ -21,6 +21,7 @@ ...@@ -21,6 +21,7 @@
#include <utility> #include <utility>
#include "caf/logger.hpp" #include "caf/logger.hpp"
#include "caf/ref_counted.hpp"
#include "caf/string_algorithms.hpp" #include "caf/string_algorithms.hpp"
#include "caf/opencl/global.hpp" #include "caf/opencl/global.hpp"
...@@ -32,71 +33,77 @@ using namespace std; ...@@ -32,71 +33,77 @@ using namespace std;
namespace caf { namespace caf {
namespace opencl { namespace opencl {
device device::create(const context_ptr& context, const device_ptr& device_id, device_ptr device::create(const cl_context_ptr& context,
unsigned id) { const cl_device_ptr& device_id,
unsigned id) {
CAF_LOG_DEBUG("creating device for opencl device with id:" << CAF_ARG(id)); CAF_LOG_DEBUG("creating device for opencl device with id:" << CAF_ARG(id));
// look up properties we need to create the command queue // look up properties we need to create the command queue
auto supported = info<cl_ulong>(device_id, CL_DEVICE_QUEUE_PROPERTIES); auto supported = info<cl_ulong>(device_id, CL_DEVICE_QUEUE_PROPERTIES);
bool profiling = (supported & CL_QUEUE_PROFILING_ENABLE) != 0u; bool profiling = false; // (supported & CL_QUEUE_PROFILING_ENABLE) != 0u;
bool out_of_order = (supported & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) != 0u; bool out_of_order = (supported & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) != 0u;
unsigned properties = profiling ? CL_QUEUE_PROFILING_ENABLE : 0; unsigned properties = profiling ? CL_QUEUE_PROFILING_ENABLE : 0;
properties |= out_of_order ? CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE : 0; properties |= out_of_order ? CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE : 0;
// create the command queue // create the command queue
command_queue_ptr command_queue; cl_command_queue_ptr command_queue{v2get(CAF_CLF(clCreateCommandQueue),
command_queue.reset(v2get(CAF_CLF(clCreateCommandQueue), context.get(), context.get(), device_id.get(),
device_id.get(), properties), properties),
false); false};
// create the device // create the device
device dev{device_id, command_queue, context, id}; auto dev = make_counted<device>(device_id, std::move(command_queue), context, id);
//device dev{device_id, std::move(command_queue), context, id};
// look up device properties // look up device properties
dev.address_bits_ = info<cl_uint>(device_id, CL_DEVICE_ADDRESS_BITS); dev->address_bits_ = info<cl_uint>(device_id, CL_DEVICE_ADDRESS_BITS);
dev.little_endian_ = info<cl_bool>(device_id, CL_DEVICE_ENDIAN_LITTLE); dev->little_endian_ = info<cl_bool>(device_id, CL_DEVICE_ENDIAN_LITTLE);
dev.global_mem_cache_size_ = dev->global_mem_cache_size_ =
info<cl_ulong>(device_id, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE); info<cl_ulong>(device_id, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE);
dev.global_mem_cacheline_size_ = dev->global_mem_cacheline_size_ =
info<cl_uint>(device_id, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE); info<cl_uint>(device_id, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE);
dev.global_mem_size_ = dev->global_mem_size_ =
info<cl_ulong >(device_id,CL_DEVICE_GLOBAL_MEM_SIZE); info<cl_ulong >(device_id,CL_DEVICE_GLOBAL_MEM_SIZE);
dev.host_unified_memory_ = dev->host_unified_memory_ =
info<cl_bool>(device_id, CL_DEVICE_HOST_UNIFIED_MEMORY); info<cl_bool>(device_id, CL_DEVICE_HOST_UNIFIED_MEMORY);
dev.local_mem_size_ = dev->local_mem_size_ =
info<cl_ulong>(device_id, CL_DEVICE_LOCAL_MEM_SIZE); info<cl_ulong>(device_id, CL_DEVICE_LOCAL_MEM_SIZE);
dev.local_mem_type_ = dev->local_mem_type_ =
info<cl_uint>(device_id, CL_DEVICE_LOCAL_MEM_TYPE); info<cl_uint>(device_id, CL_DEVICE_LOCAL_MEM_TYPE);
dev.max_clock_frequency_ = dev->max_clock_frequency_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_CLOCK_FREQUENCY); info<cl_uint>(device_id, CL_DEVICE_MAX_CLOCK_FREQUENCY);
dev.max_compute_units_ = dev->max_compute_units_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_COMPUTE_UNITS); info<cl_uint>(device_id, CL_DEVICE_MAX_COMPUTE_UNITS);
dev.max_constant_args_ = dev->max_constant_args_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_CONSTANT_ARGS); info<cl_uint>(device_id, CL_DEVICE_MAX_CONSTANT_ARGS);
dev.max_constant_buffer_size_ = dev->max_constant_buffer_size_ =
info<cl_ulong>(device_id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE); info<cl_ulong>(device_id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE);
dev.max_mem_alloc_size_ = dev->max_mem_alloc_size_ =
info<cl_ulong>(device_id, CL_DEVICE_MAX_MEM_ALLOC_SIZE); info<cl_ulong>(device_id, CL_DEVICE_MAX_MEM_ALLOC_SIZE);
dev.max_parameter_size_ = dev->max_parameter_size_ =
info<size_t>(device_id, CL_DEVICE_MAX_PARAMETER_SIZE); info<size_t>(device_id, CL_DEVICE_MAX_PARAMETER_SIZE);
dev.max_work_group_size_ = dev->max_work_group_size_ =
info<size_t>(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE); info<size_t>(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE);
dev.max_work_item_dimensions_ = dev->max_work_item_dimensions_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS); info<cl_uint>(device_id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS);
dev.profiling_timer_resolution_ = dev->profiling_timer_resolution_ =
info<size_t>(device_id, CL_DEVICE_PROFILING_TIMER_RESOLUTION); info<size_t>(device_id, CL_DEVICE_PROFILING_TIMER_RESOLUTION);
dev.max_work_item_sizes_.resize(dev.max_work_item_dimensions_); dev->max_work_item_sizes_.resize(dev->max_work_item_dimensions_);
clGetDeviceInfo(device_id.get(), CL_DEVICE_MAX_WORK_ITEM_SIZES, clGetDeviceInfo(device_id.get(), CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(size_t) * dev.max_work_item_dimensions_, sizeof(size_t) * dev->max_work_item_dimensions_,
dev.max_work_item_sizes_.data(), nullptr); dev->max_work_item_sizes_.data(), nullptr);
dev.device_type_ = dev->device_type_ =
device_type_from_ulong(info<cl_ulong>(device_id, CL_DEVICE_TYPE)); device_type_from_ulong(info<cl_ulong>(device_id, CL_DEVICE_TYPE));
string extensions = info_string(device_id, CL_DEVICE_EXTENSIONS); string extensions = info_string(device_id, CL_DEVICE_EXTENSIONS);
split(dev.extensions_, extensions, " ", false); split(dev->extensions_, extensions, " ", false);
dev.opencl_c_version_ = info_string(device_id, CL_DEVICE_EXTENSIONS); dev->opencl_c_version_ = info_string(device_id, CL_DEVICE_EXTENSIONS);
dev.device_vendor_ = info_string(device_id, CL_DEVICE_VENDOR); dev->device_vendor_ = info_string(device_id, CL_DEVICE_VENDOR);
dev.device_version_ = info_string(device_id, CL_DEVICE_VERSION); dev->device_version_ = info_string(device_id, CL_DEVICE_VERSION);
dev.name_ = info_string(device_id, CL_DEVICE_NAME); dev->name_ = info_string(device_id, CL_DEVICE_NAME);
return dev; return dev;
} }
string device::info_string(const device_ptr& device_id, unsigned info_flag) { void device::synchronize() {
clFinish(queue_.get());
}
string device::info_string(const cl_device_ptr& device_id, unsigned info_flag) {
size_t size; size_t size;
clGetDeviceInfo(device_id.get(), info_flag, 0, nullptr, &size); clGetDeviceInfo(device_id.get(), info_flag, 0, nullptr, &size);
vector<char> buffer(size); vector<char> buffer(size);
...@@ -105,12 +112,18 @@ string device::info_string(const device_ptr& device_id, unsigned info_flag) { ...@@ -105,12 +112,18 @@ string device::info_string(const device_ptr& device_id, unsigned info_flag) {
return string(buffer.data()); return string(buffer.data());
} }
device::device(device_ptr device_id, command_queue_ptr queue, device::device(cl_device_ptr device_id, cl_command_queue_ptr queue,
context_ptr context, unsigned id) cl_context_ptr context, unsigned id)
: device_id_(std::move(device_id)), : device_id_(std::move(device_id)),
command_queue_(std::move(queue)), queue_(std::move(queue)),
context_(std::move(context)), context_(std::move(context)),
id_(id) { } id_(id) {
// nop
}
device::~device() {
// nop
}
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
...@@ -18,6 +18,7 @@ ...@@ -18,6 +18,7 @@
******************************************************************************/ ******************************************************************************/
#include <string> #include <string>
#include <sstream>
#include "caf/opencl/global.hpp" #include "caf/opencl/global.hpp"
...@@ -171,5 +172,122 @@ std::string get_opencl_error(cl_int err) { ...@@ -171,5 +172,122 @@ std::string get_opencl_error(cl_int err) {
} }
} }
std::string event_status(cl_event e) {
std::stringstream ss;
cl_int s;
auto err = clGetEventInfo(e, CL_EVENT_COMMAND_EXECUTION_STATUS,
sizeof(cl_int), &s, nullptr);
if (err != CL_SUCCESS) {
ss << std::string("ERROR ") + std::to_string(s);
return ss.str();
}
cl_command_type t;
err = clGetEventInfo(e, CL_EVENT_COMMAND_TYPE, sizeof(cl_command_type),
&t, nullptr);
if (err != CL_SUCCESS) {
ss << std::string("ERROR ") + std::to_string(s);
return ss.str();
}
switch (s) {
case(CL_QUEUED):
ss << std::string("CL_QUEUED");
break;
case(CL_SUBMITTED):
ss << std::string("CL_SUBMITTED");
break;
case(CL_RUNNING):
ss << std::string("CL_RUNNING");
break;
case(CL_COMPLETE):
ss << std::string("CL_COMPLETE");
break;
default:
ss << std::string("DEFAULT ") + std::to_string(s);
return ss.str();
}
ss << " / ";
switch (t) {
case(CL_COMMAND_NDRANGE_KERNEL):
ss << std::string("CL_COMMAND_NDRANGE_KERNEL");
break;
case(CL_COMMAND_TASK):
ss << std::string("CL_COMMAND_TASK");
break;
case(CL_COMMAND_NATIVE_KERNEL):
ss << std::string("CL_COMMAND_NATIVE_KERNEL");
break;
case(CL_COMMAND_READ_BUFFER):
ss << std::string("CL_COMMAND_READ_BUFFER");
break;
case(CL_COMMAND_WRITE_BUFFER):
ss << std::string("CL_COMMAND_WRITE_BUFFER");
break;
case(CL_COMMAND_COPY_BUFFER):
ss << std::string("CL_COMMAND_COPY_BUFFER");
break;
case(CL_COMMAND_READ_IMAGE):
ss << std::string("CL_COMMAND_READ_IMAGE");
break;
case(CL_COMMAND_WRITE_IMAGE):
ss << std::string("CL_COMMAND_WRITE_IMAGE");
break;
case(CL_COMMAND_COPY_IMAGE):
ss << std::string("CL_COMMAND_COPY_IMAGE");
break;
case(CL_COMMAND_COPY_BUFFER_TO_IMAGE):
ss << std::string("CL_COMMAND_COPY_BUFFER_TO_IMAGE");
break;
case(CL_COMMAND_COPY_IMAGE_TO_BUFFER):
ss << std::string("CL_COMMAND_COPY_IMAGE_TO_BUFFER");
break;
case(CL_COMMAND_MAP_BUFFER):
ss << std::string("CL_COMMAND_MAP_BUFFER");
break;
case(CL_COMMAND_MAP_IMAGE):
ss << std::string("CL_COMMAND_MAP_IMAGE");
break;
case(CL_COMMAND_UNMAP_MEM_OBJECT):
ss << std::string("CL_COMMAND_UNMAP_MEM_OBJECT");
break;
case(CL_COMMAND_MARKER):
ss << std::string("CL_COMMAND_MARKER");
break;
case(CL_COMMAND_ACQUIRE_GL_OBJECTS):
ss << std::string("CL_COMMAND_ACQUIRE_GL_OBJECTS");
break;
case(CL_COMMAND_RELEASE_GL_OBJECTS):
ss << std::string("CL_COMMAND_RELEASE_GL_OBJECTS");
break;
case(CL_COMMAND_READ_BUFFER_RECT):
ss << std::string("CL_COMMAND_READ_BUFFER_RECT");
break;
case(CL_COMMAND_WRITE_BUFFER_RECT):
ss << std::string("CL_COMMAND_WRITE_BUFFER_RECT");
break;
case(CL_COMMAND_COPY_BUFFER_RECT):
ss << std::string("CL_COMMAND_COPY_BUFFER_RECT");
break;
case(CL_COMMAND_USER):
ss << std::string("CL_COMMAND_USER");
break;
case(CL_COMMAND_BARRIER):
ss << std::string("CL_COMMAND_BARRIER");
break;
case(CL_COMMAND_MIGRATE_MEM_OBJECTS):
ss << std::string("CL_COMMAND_MIGRATE_MEM_OBJECTS");
break;
case(CL_COMMAND_FILL_BUFFER):
ss << std::string("CL_COMMAND_FILL_BUFFER");
break;
case(CL_COMMAND_FILL_IMAGE):
ss << std::string("CL_COMMAND_FILL_IMAGE");
break;
default:
ss << std::string("DEFAULT ") + std::to_string(s);
return ss.str();
}
return ss.str();
}
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
...@@ -18,11 +18,14 @@ ...@@ -18,11 +18,14 @@
* http://www.boost.org/LICENSE_1_0.txt. * * http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/ ******************************************************************************/
#include <fstream>
#include "caf/detail/type_list.hpp" #include "caf/detail/type_list.hpp"
#include "caf/opencl/device.hpp" #include "caf/opencl/device.hpp"
#include "caf/opencl/manager.hpp" #include "caf/opencl/manager.hpp"
#include "caf/opencl/platform.hpp" #include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp" #include "caf/opencl/opencl_err.hpp"
using namespace std; using namespace std;
...@@ -30,15 +33,15 @@ using namespace std; ...@@ -30,15 +33,15 @@ using namespace std;
namespace caf { namespace caf {
namespace opencl { namespace opencl {
const optional<const device&> manager::get_device(size_t dev_id) const { optional<device_ptr> manager::get_device(size_t dev_id) const {
if (platforms_.empty()) if (platforms_.empty())
return none; return none;
size_t to = 0; size_t to = 0;
for (auto& pl : platforms_) { for (auto& pl : platforms_) {
auto from = to; auto from = to;
to += pl.get_devices().size(); to += pl->get_devices().size();
if (dev_id >= from && dev_id < to) if (dev_id >= from && dev_id < to)
return pl.get_devices()[dev_id - from]; return pl->get_devices()[dev_id - from];
} }
return none; return none;
} }
...@@ -56,7 +59,7 @@ void manager::init(actor_system_config&) { ...@@ -56,7 +59,7 @@ void manager::init(actor_system_config&) {
for (auto& pl_id : platform_ids) { for (auto& pl_id : platform_ids) {
platforms_.push_back(platform::create(pl_id, current_device_id)); platforms_.push_back(platform::create(pl_id, current_device_id));
current_device_id += current_device_id +=
static_cast<unsigned>(platforms_.back().get_devices().size()); static_cast<unsigned>(platforms_.back()->get_devices().size());
} }
} }
...@@ -81,8 +84,30 @@ actor_system::module* manager::make(actor_system& sys, ...@@ -81,8 +84,30 @@ actor_system::module* manager::make(actor_system& sys,
return new manager{sys}; return new manager{sys};
} }
program manager::create_program(const char* kernel_source, const char* options, program_ptr manager::create_program_from_file(const char* path,
uint32_t device_id) { const char* options,
uint32_t device_id) {
std::ifstream read_source{std::string(path), std::ios::in};
string kernel_source;
if (read_source) {
read_source.seekg(0, std::ios::end);
kernel_source.resize(static_cast<size_t>(read_source.tellg()));
read_source.seekg(0, std::ios::beg);
read_source.read(&kernel_source[0],
static_cast<streamsize>(kernel_source.size()));
read_source.close();
} else {
ostringstream oss;
oss << "No file at '" << path << "' found.";
CAF_LOG_ERROR(CAF_ARG(oss.str()));
throw runtime_error(oss.str());
}
return create_program(kernel_source.c_str(), options, device_id);
}
program_ptr manager::create_program(const char* kernel_source,
const char* options,
uint32_t device_id) {
auto dev = get_device(device_id); auto dev = get_device(device_id);
if (!dev) { if (!dev) {
ostringstream oss; ostringstream oss;
...@@ -93,61 +118,83 @@ program manager::create_program(const char* kernel_source, const char* options, ...@@ -93,61 +118,83 @@ program manager::create_program(const char* kernel_source, const char* options,
return create_program(kernel_source, options, *dev); return create_program(kernel_source, options, *dev);
} }
program manager::create_program(const char* kernel_source, const char* options, program_ptr manager::create_program_from_file(const char* path,
const device& dev) { const char* options,
const device_ptr dev) {
std::ifstream read_source{std::string(path), std::ios::in};
string kernel_source;
if (read_source) {
read_source.seekg(0, std::ios::end);
kernel_source.resize(static_cast<size_t>(read_source.tellg()));
read_source.seekg(0, std::ios::beg);
read_source.read(&kernel_source[0],
static_cast<streamsize>(kernel_source.size()));
read_source.close();
} else {
ostringstream oss;
oss << "No file at '" << path << "' found.";
CAF_LOG_ERROR(CAF_ARG(oss.str()));
throw runtime_error(oss.str());
}
return create_program(kernel_source.c_str(), options, dev);
}
program_ptr manager::create_program(const char* kernel_source,
const char* options,
const device_ptr dev) {
// create program object from kernel source // create program object from kernel source
size_t kernel_source_length = strlen(kernel_source); size_t kernel_source_length = strlen(kernel_source);
program_ptr pptr; cl_program_ptr pptr;
auto rawptr = v2get(CAF_CLF(clCreateProgramWithSource), dev.context_.get(), pptr.reset(v2get(CAF_CLF(clCreateProgramWithSource), dev->context_.get(),
cl_uint{1}, &kernel_source, &kernel_source_length); 1u, &kernel_source, &kernel_source_length),
pptr.reset(rawptr, false); false);
// build programm from program object // build programm from program object
auto dev_tmp = dev.device_id_.get(); auto dev_tmp = dev->device_id_.get();
cl_int err = clBuildProgram(pptr.get(), 1, &dev_tmp, auto err = clBuildProgram(pptr.get(), 1, &dev_tmp, options, nullptr, nullptr);
options, nullptr, nullptr);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
ostringstream oss; ostringstream oss;
oss << "clBuildProgram: " << get_opencl_error(err); oss << "clBuildProgram: " << get_opencl_error(err);
// the build log will be printed by the pfn_notify (see opencl/manger.cpp)
#ifndef __APPLE__
// seems that just apple implemented the
// pfn_notify callback, but we can get
// the build log
if (err == CL_BUILD_PROGRAM_FAILURE) { if (err == CL_BUILD_PROGRAM_FAILURE) {
size_t buildlog_buffer_size = 0; size_t buildlog_buffer_size = 0;
// get the log length // get the log length
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG, clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buildlog_buffer_size), nullptr, 0, nullptr, &buildlog_buffer_size);
&buildlog_buffer_size);
vector<char> buffer(buildlog_buffer_size); vector<char> buffer(buildlog_buffer_size);
// fill the buffer with buildlog informations // fill the buffer with buildlog informations
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG, clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buffer[0]) * buildlog_buffer_size, sizeof(char) * buildlog_buffer_size,
buffer.data(), nullptr); buffer.data(), nullptr);
ostringstream ss; ostringstream ss;
ss << "Build log:\n" << string(buffer.data()) ss << "############## Build log ##############"
<< "\n#######################################"; << endl << string(buffer.data()) << endl
<< "#######################################";
// seems that just apple implemented the
// pfn_notify callback, but we can get
// the build log
#ifndef __APPLE__
CAF_LOG_ERROR(CAF_ARG(ss.str())); CAF_LOG_ERROR(CAF_ARG(ss.str()));
}
#endif #endif
oss << endl << ss.str();
}
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
map<string, kernel_ptr> available_kernels;
cl_uint number_of_kernels = 0; cl_uint number_of_kernels = 0;
clCreateKernelsInProgram(pptr.get(), 0, nullptr, &number_of_kernels); clCreateKernelsInProgram(pptr.get(), 0u, nullptr, &number_of_kernels);
vector<cl_kernel> kernels(number_of_kernels); map<string, cl_kernel_ptr> available_kernels;
err = clCreateKernelsInProgram(pptr.get(), number_of_kernels, kernels.data(), if (number_of_kernels > 0) {
nullptr); vector<cl_kernel> kernels(number_of_kernels);
if (err != CL_SUCCESS) { err = clCreateKernelsInProgram(pptr.get(), number_of_kernels,
ostringstream oss; kernels.data(), nullptr);
oss << "clCreateKernelsInProgram: " << get_opencl_error(err); if (err != CL_SUCCESS) {
throw runtime_error(oss.str()); ostringstream oss;
} else { oss << "clCreateKernelsInProgram: " << get_opencl_error(err);
throw runtime_error(oss.str());
}
for (cl_uint i = 0; i < number_of_kernels; ++i) { for (cl_uint i = 0; i < number_of_kernels; ++i) {
size_t ret_size; size_t len;
clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, 0, nullptr, &ret_size); clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, 0, nullptr, &len);
vector<char> name(ret_size); vector<char> name(len);
err = clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, ret_size, err = clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, len,
reinterpret_cast<void*>(name.data()), nullptr); reinterpret_cast<void*>(name.data()), nullptr);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
ostringstream oss; ostringstream oss;
...@@ -155,15 +202,20 @@ program manager::create_program(const char* kernel_source, const char* options, ...@@ -155,15 +202,20 @@ program manager::create_program(const char* kernel_source, const char* options,
<< get_opencl_error(err); << get_opencl_error(err);
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
kernel_ptr kernel; cl_kernel_ptr kernel;
kernel.reset(move(kernels[i])); kernel.reset(move(kernels[i]));
available_kernels.emplace(string(name.data()), move(kernel)); available_kernels.emplace(string(name.data()), move(kernel));
} }
} else {
CAF_LOG_WARNING("Could not built all kernels in program. Since this happens"
" on some platforms, we'll ignore this and try to build"
" each kernel individually by name.");
} }
return {dev.context_, dev.command_queue_, pptr, move(available_kernels)}; return make_counted<program>(dev->context_, dev->queue_, pptr,
move(available_kernels));
} }
manager::manager(actor_system& sys) : system_(sys){ manager::manager(actor_system& sys) : system_(sys) {
// nop // nop
} }
......
...@@ -29,8 +29,8 @@ using namespace std; ...@@ -29,8 +29,8 @@ using namespace std;
namespace caf { namespace caf {
namespace opencl { namespace opencl {
platform platform::create(cl_platform_id platform_id, platform_ptr platform::create(cl_platform_id platform_id,
unsigned start_id) { unsigned start_id) {
vector<unsigned> device_types = {CL_DEVICE_TYPE_GPU, vector<unsigned> device_types = {CL_DEVICE_TYPE_GPU,
CL_DEVICE_TYPE_ACCELERATOR, CL_DEVICE_TYPE_ACCELERATOR,
CL_DEVICE_TYPE_CPU}; CL_DEVICE_TYPE_CPU};
...@@ -48,16 +48,16 @@ platform platform::create(cl_platform_id platform_id, ...@@ -48,16 +48,16 @@ platform platform::create(cl_platform_id platform_id,
v2callcl(CAF_CLF(clGetDeviceIDs), platform_id, device_type, v2callcl(CAF_CLF(clGetDeviceIDs), platform_id, device_type,
discoverd, (ids.data() + known)); discoverd, (ids.data() + known));
} }
vector<device_ptr> devices; vector<cl_device_ptr> devices;
devices.resize(ids.size()); devices.resize(ids.size());
auto lift = [](cl_device_id ptr) { return device_ptr{ptr, false}; }; auto lift = [](cl_device_id ptr) { return cl_device_ptr{ptr, false}; };
transform(ids.begin(), ids.end(), devices.begin(), lift); transform(ids.begin(), ids.end(), devices.begin(), lift);
context_ptr context; cl_context_ptr context;
context.reset(v2get(CAF_CLF(clCreateContext), nullptr, context.reset(v2get(CAF_CLF(clCreateContext), nullptr,
static_cast<unsigned>(ids.size()), static_cast<unsigned>(ids.size()),
ids.data(), pfn_notify, nullptr), ids.data(), pfn_notify, nullptr),
false); false);
vector<device> device_information; vector<device_ptr> device_information;
for (auto& device_id : devices) { for (auto& device_id : devices) {
device_information.push_back(device::create(context, device_id, device_information.push_back(device::create(context, device_id,
start_id++)); start_id++));
...@@ -70,9 +70,9 @@ platform platform::create(cl_platform_id platform_id, ...@@ -70,9 +70,9 @@ platform platform::create(cl_platform_id platform_id,
auto name = platform_info(platform_id, CL_PLATFORM_NAME); auto name = platform_info(platform_id, CL_PLATFORM_NAME);
auto vendor = platform_info(platform_id, CL_PLATFORM_VENDOR); auto vendor = platform_info(platform_id, CL_PLATFORM_VENDOR);
auto version = platform_info(platform_id, CL_PLATFORM_VERSION); auto version = platform_info(platform_id, CL_PLATFORM_VERSION);
return platform{platform_id, move(context), move(name), return make_counted<platform>(platform_id, move(context), move(name),
move(vendor), move(version), move(vendor), move(version),
move(device_information)}; move(device_information));
} }
string platform::platform_info(cl_platform_id platform_id, string platform::platform_info(cl_platform_id platform_id,
...@@ -87,15 +87,21 @@ string platform::platform_info(cl_platform_id platform_id, ...@@ -87,15 +87,21 @@ string platform::platform_info(cl_platform_id platform_id,
return string(buffer.data()); return string(buffer.data());
} }
platform::platform(cl_platform_id platform_id, context_ptr context, platform::platform(cl_platform_id platform_id, cl_context_ptr context,
string name, string vendor, string version, string name, string vendor, string version,
vector<device> devices) vector<device_ptr> devices)
: platform_id_(platform_id), : platform_id_(platform_id),
context_(std::move(context)), context_(std::move(context)),
name_(move(name)), name_(move(name)),
vendor_(move(vendor)), vendor_(move(vendor)),
version_(move(version)), version_(move(version)),
devices_(move(devices)) { } devices_(move(devices)) {
// nop
}
platform::~platform() {
// nop
}
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
...@@ -32,8 +32,9 @@ using namespace std; ...@@ -32,8 +32,9 @@ using namespace std;
namespace caf { namespace caf {
namespace opencl { namespace opencl {
program::program(context_ptr context, command_queue_ptr queue, program_ptr prog, program::program(cl_context_ptr context, cl_command_queue_ptr queue,
map<string, kernel_ptr> available_kernels) cl_program_ptr prog,
map<string, cl_kernel_ptr> available_kernels)
: context_(move(context)), : context_(move(context)),
program_(move(prog)), program_(move(prog)),
queue_(move(queue)), queue_(move(queue)),
...@@ -41,5 +42,9 @@ program::program(context_ptr context, command_queue_ptr queue, program_ptr prog, ...@@ -41,5 +42,9 @@ program::program(context_ptr context, command_queue_ptr queue, program_ptr prog,
// nop // nop
} }
program::~program() {
// nop
}
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
...@@ -18,26 +18,35 @@ using namespace caf::opencl; ...@@ -18,26 +18,35 @@ using namespace caf::opencl;
using caf::detail::limited_vector; using caf::detail::limited_vector;
// required to allow sending mem_ref<int> in messages
CAF_ALLOW_UNSAFE_MESSAGE_TYPE(mem_ref<int>);
namespace { namespace {
using ivec = vector<int>; using ivec = vector<int>;
using iref = mem_ref<int>;
using dims = opencl::dim_vec; using dims = opencl::dim_vec;
constexpr size_t matrix_size = 4; constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32; constexpr size_t array_size = 32;
constexpr size_t problem_size = 1024; constexpr size_t problem_size = 1024;
constexpr const char* kernel_name = "matrix_square"; constexpr const char* kn_matrix = "matrix_square";
constexpr const char* kernel_name_compiler_flag = "compiler_flag"; constexpr const char* kn_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_reduce = "reduce"; constexpr const char* kn_reduce = "reduce";
constexpr const char* kernel_name_const = "const_mod"; constexpr const char* kn_const = "const_mod";
constexpr const char* kernel_name_inout = "times_two"; constexpr const char* kn_inout = "times_two";
constexpr const char* kn_scratch = "use_scratch";
constexpr const char* kn_local = "use_local";
constexpr const char* kn_order = "test_order";
constexpr const char* kn_private = "use_private";
constexpr const char* kn_varying = "varying";
constexpr const char* compiler_flag = "-D CAF_OPENCL_TEST_FLAG"; constexpr const char* compiler_flag = "-D CAF_OPENCL_TEST_FLAG";
constexpr const char* kernel_source = R"__( constexpr const char* kernel_source = R"__(
__kernel void matrix_square(__global int* matrix, kernel void matrix_square(global const int* restrict matrix,
__global int* output) { global int* restrict output) {
size_t size = get_global_size(0); // == get_global_size_(1); size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0); size_t x = get_global_id(0);
size_t y = get_global_id(1); size_t y = get_global_id(1);
...@@ -47,32 +56,12 @@ constexpr const char* kernel_source = R"__( ...@@ -47,32 +56,12 @@ constexpr const char* kernel_source = R"__(
} }
output[x + y * size] = result; output[x + y * size] = result;
} }
)__";
constexpr const char* kernel_source_error = R"__(
__kernel void missing(__global int*) {
size_t semicolon_missing
}
)__";
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/ // http://developer.amd.com/resources/documentation-articles/
// opencl-optimization-case-study-simple-reductions // articles-whitepapers/opencl-optimization-case-study-simple-reductions
constexpr const char* kernel_source_reduce = R"__( kernel void reduce(global const int* restrict buffer,
__kernel void reduce(__global int* buffer, global int* restrict result) {
__global int* result) { local int scratch[512];
__local int scratch[512];
int local_index = get_local_id(0); int local_index = get_local_id(0);
scratch[local_index] = buffer[get_global_id(0)]; scratch[local_index] = buffer[get_global_id(0)];
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
...@@ -84,25 +73,110 @@ constexpr const char* kernel_source_reduce = R"__( ...@@ -84,25 +73,110 @@ constexpr const char* kernel_source_reduce = R"__(
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
} }
if (local_index == 0) { if (local_index == 0)
result[get_group_id(0)] = scratch[0]; result[get_group_id(0)] = scratch[0];
}
} }
)__";
constexpr const char* kernel_source_const = R"__( kernel void const_mod(constant int* restrict input,
__kernel void const_mod(__constant int* input, global int* restrict output) {
__global int* output) {
size_t idx = get_global_id(0); size_t idx = get_global_id(0);
output[idx] = input[0]; output[idx] = input[0];
} }
)__";
constexpr const char* kernel_source_inout = R"__( kernel void times_two(global int* restrict values) {
__kernel void times_two(__global int* values) {
size_t idx = get_global_id(0); size_t idx = get_global_id(0);
values[idx] = values[idx] * 2; values[idx] = values[idx] * 2;
} }
kernel void use_scratch(global int* restrict values,
global int* restrict buf) {
size_t idx = get_global_id(0);
buf[idx] = values[idx];
buf[idx] += values[idx];
values[idx] = buf[idx];
}
inline void prefix_sum(local int* restrict data, size_t len, size_t lids) {
size_t lid = get_local_id(0);
size_t inc = 2;
// reduce
while (inc <= len) {
int j = inc >> 1;
for (int i = (j - 1) + (lid * inc); (i + inc) < len; i += (lids * inc))
data[i + j] = data[i] + data[i + j];
inc = inc << 1;
barrier(CLK_LOCAL_MEM_FENCE);
}
// downsweep
data[len - 1] = 0;
barrier(CLK_LOCAL_MEM_FENCE);
while (inc >= 2) {
int j = inc >> 1;
for (int i = (j - 1) + (lid * inc); (i + j) <= len; i += (lids * inc)) {
uint tmp = data[i + j];
data[i + j] = data[i] + data[i + j];
data[i] = tmp;
}
inc = inc >> 1;
barrier(CLK_LOCAL_MEM_FENCE);
}
}
kernel void use_local(global int* restrict values,
local int* restrict buf) {
size_t lid = get_local_id(0);
size_t gid = get_group_id(0);
size_t gs = get_local_size(0);
buf[lid] = values[gid * gs + lid];
barrier(CLK_LOCAL_MEM_FENCE);
prefix_sum(buf, gs, gs);
barrier(CLK_LOCAL_MEM_FENCE);
values[gid * gs + lid] = buf[lid];
}
kernel void test_order(local int* buf,
global int* restrict values) {
size_t lid = get_local_id(0);
size_t gid = get_group_id(0);
size_t gs = get_local_size(0);
buf[lid] = values[gid * gs + lid];
barrier(CLK_LOCAL_MEM_FENCE);
prefix_sum(buf, gs, gs);
barrier(CLK_LOCAL_MEM_FENCE);
values[gid * gs + lid] = buf[lid];
}
kernel void use_private(global int* restrict buf,
private int val) {
buf[get_global_id(0)] += val;
}
kernel void varying(global const int* restrict in1,
global int* restrict out1,
global const int* restrict in2,
global int* restrict out2) {
size_t idx = get_global_id(0);
out1[idx] = in1[idx];
out2[idx] = in2[idx];
}
)__";
constexpr const char* kernel_source_error = R"__(
kernel void missing(global int*) {
size_t semicolon_missing
}
)__";
constexpr const char* kernel_source_compiler_flag = R"__(
kernel void compiler_flag(global const int* restrict input,
global int* restrict output) {
size_t x = get_global_id(0);
# ifdef CAF_OPENCL_TEST_FLAG
output[x] = input[x];
# else
output[x] = 0;
# endif
}
)__"; )__";
} // namespace <anonymous> } // namespace <anonymous>
...@@ -213,49 +287,81 @@ void check_vector_results(const string& description, ...@@ -213,49 +287,81 @@ void check_vector_results(const string& description,
cout << " " << result[i]; cout << " " << result[i];
} }
cout << endl; cout << endl;
cout << "Size: " << expected.size() << " vs. " << result.size() << endl;
cout << "Differ at: " << endl;
bool same = true;
for (size_t i = 0; i < min(expected.size(), result.size()); ++i) {
if (expected[i] != result[i]) {
cout << "[" << i << "] " << expected[i] << " != " << result[i] << endl;
same = false;
}
}
if (same) {
cout << "... nowhere." << endl;
}
}
}
template <class T>
void check_mref_results(const string& description,
const vector<T>& expected,
mem_ref<T>& result) {
auto exp_res = result.data();
CAF_REQUIRE(exp_res);
auto res = *exp_res;
auto cond = (expected == res);
CAF_CHECK(cond);
if (!cond) {
CAF_ERROR(description << " failed.");
cout << "Expected: " << endl;
for (size_t i = 0; i < expected.size(); ++i) {
cout << " " << expected[i];
}
cout << endl << "Received: " << endl;
for (size_t i = 0; i < res.size(); ++i) {
cout << " " << res[i];
}
cout << endl;
} }
} }
void test_opencl(actor_system& sys) { void test_opencl(actor_system& sys) {
static_cast<void>(sys);
auto& mngr = sys.opencl_manager(); auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device_if([](const device&){ return true; }); auto opt = mngr.get_device_if([](const device_ptr){ return true; });
if (!opt) CAF_REQUIRE(opt);
CAF_ERROR("No OpenCL device found.");
auto dev = *opt; auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys}; scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
const ivec expected1{ 56, 62, 68, 74, const ivec expected1{ 56, 62, 68, 74,
152, 174, 196, 218, 152, 174, 196, 218,
248, 286, 324, 362, 248, 286, 324, 362,
344, 398, 452, 506}; 344, 398, 452, 506};
auto w1 = mngr.spawn(mngr.create_program(kernel_source, "", dev), kernel_name, auto w1 = mngr.spawn(prog, kn_matrix,
opencl::spawn_config{dims{matrix_size, matrix_size}}, opencl::spawn_config{dims{matrix_size, matrix_size}},
opencl::in<int*>{}, opencl::out<int*>{}); opencl::in<int>{}, opencl::out<int>{});
self->send(w1, make_iota_vector<int>(matrix_size * matrix_size)); self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
self->receive ( self->receive (
[&](const ivec& result) { [&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors" check_vector_results("Simple matrix multiplication using vectors"
"(kernel wrapped in program)", " (kernel wrapped in program)",
expected1, result); expected1, result);
}, }, others >> wrong_msg
others >> [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
}
); );
opencl::spawn_config cfg2{dims{matrix_size, matrix_size}}; opencl::spawn_config cfg2{dims{matrix_size, matrix_size}};
auto w2 = mngr.spawn(kernel_source, kernel_name, cfg2, // Pass kernel directly to the actor
opencl::in<int*>{}, opencl::out<int*>{}); auto w2 = mngr.spawn(kernel_source, kn_matrix, cfg2,
opencl::in<int>{}, opencl::out<int>{});
self->send(w2, make_iota_vector<int>(matrix_size * matrix_size)); self->send(w2, make_iota_vector<int>(matrix_size * matrix_size));
self->receive ( self->receive (
[&](const ivec& result) { [&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors", check_vector_results("Simple matrix multiplication using vectors"
" (kernel passed directly)",
expected1, result); expected1, result);
}, }, others >> wrong_msg
others >> [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
}
); );
const matrix_type expected2(move(expected1)); const matrix_type expected2(move(expected1));
auto map_arg = [](message& msg) -> optional<message> { auto map_arg = [](message& msg) -> optional<message> {
...@@ -265,151 +371,209 @@ void test_opencl(actor_system& sys) { ...@@ -265,151 +371,209 @@ void test_opencl(actor_system& sys) {
} }
); );
}; };
auto map_res = [=](ivec result) -> message { auto map_res = [](ivec result) -> message {
return make_message(matrix_type{move(result)}); return make_message(matrix_type{move(result)});
}; };
opencl::spawn_config cfg3{dims{matrix_size, matrix_size}}; opencl::spawn_config cfg3{dims{matrix_size, matrix_size}};
auto w3 = mngr.spawn(mngr.create_program(kernel_source), kernel_name, cfg3, // let the runtime choose the device
auto w3 = mngr.spawn(mngr.create_program(kernel_source), kn_matrix, cfg3,
map_arg, map_res, map_arg, map_res,
opencl::in<int*>{}, opencl::out<int*>{}); opencl::in<int>{}, opencl::out<int>{});
self->send(w3, make_iota_matrix<matrix_size>()); self->send(w3, make_iota_matrix<matrix_size>());
self->receive ( self->receive (
[&](const matrix_type& result) { [&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type " check_vector_results("Matrix multiplication with user defined type "
"(kernel wrapped in program)", "(kernel wrapped in program)",
expected2.data(), result.data()); expected2.data(), result.data());
}, }, others >> wrong_msg
others >> [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
}
); );
opencl::spawn_config cfg4{dims{matrix_size, matrix_size}}; opencl::spawn_config cfg4{dims{matrix_size, matrix_size}};
auto w4 = mngr.spawn(kernel_source, kernel_name, cfg4, auto w4 = mngr.spawn(prog, kn_matrix, cfg4,
map_arg, map_res, map_arg, map_res,
opencl::in<int*>{}, opencl::out<int*>{}); opencl::in<int>{}, opencl::out<int>{});
self->send(w4, make_iota_matrix<matrix_size>()); self->send(w4, make_iota_matrix<matrix_size>());
self->receive ( self->receive (
[&](const matrix_type& result) { [&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type", check_vector_results("Matrix multiplication with user defined type",
expected2.data(), result.data()); expected2.data(), result.data());
}, }, others >> wrong_msg
others >> [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
}
); );
CAF_MESSAGE("Expecting exception (compiling invalid kernel, " CAF_MESSAGE("Expecting exception (compiling invalid kernel, "
"semicolon is missing)."); "semicolon is missing).");
try { try {
auto create_error = mngr.create_program(kernel_source_error); /* auto expected_error = */ mngr.create_program(kernel_source_error);
} } catch (const exception& exc) {
catch (const exception& exc) { std::string starts_with("clBuildProgram: CL_BUILD_PROGRAM_FAILURE");
auto cond = (strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", auto cond = (strncmp(exc.what(), starts_with.c_str(),
exc.what()) == 0); starts_with.size()) == 0);
CAF_CHECK(cond); CAF_CHECK(cond);
if (!cond) { if (!cond)
CAF_ERROR("Wrong exception cought for program build failure."); CAF_ERROR("Wrong exception cought for program build failure.");
}
} }
// test for opencl compiler flags
// create program with opencl compiler flags
auto prog5 = mngr.create_program(kernel_source_compiler_flag, compiler_flag); auto prog5 = mngr.create_program(kernel_source_compiler_flag, compiler_flag);
opencl::spawn_config cfg5{dims{array_size}}; opencl::spawn_config cfg5{dims{array_size}};
auto w5 = mngr.spawn(prog5, kernel_name_compiler_flag, cfg5, auto w5 = mngr.spawn(prog5, kn_compiler_flag, cfg5,
opencl::in<int*>{}, opencl::out<int*>{}); opencl::in<int>{}, opencl::out<int>{});
self->send(w5, make_iota_vector<int>(array_size)); self->send(w5, make_iota_vector<int>(array_size));
auto expected3 = make_iota_vector<int>(array_size); auto expected3 = make_iota_vector<int>(array_size);
self->receive ( self->receive (
[&](const ivec& result) { [&](const ivec& result) {
check_vector_results("Passing compiler flags", expected3, result); check_vector_results("Passing compiler flags", expected3, result);
}, }, others >> wrong_msg
others >> [&](message_view& x) -> result<message> { );
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message; // test for manuel return size selection (max workgroup size 1d)
} auto max_wg_size = min(dev->get_max_work_item_sizes()[0], size_t{512});
auto reduce_buffer_size = static_cast<size_t>(max_wg_size) * 8;
auto reduce_local_size = static_cast<size_t>(max_wg_size);
auto reduce_work_groups = reduce_buffer_size / reduce_local_size;
auto reduce_global_size = reduce_buffer_size;
auto reduce_result_size = reduce_work_groups;
ivec arr6(reduce_buffer_size);
int n = static_cast<int>(arr6.capacity());
generate(arr6.begin(), arr6.end(), [&]{ return --n; });
opencl::spawn_config cfg6{dims{reduce_global_size},
dims{ }, // no offset
dims{reduce_local_size}};
auto result_size_6 = [reduce_result_size](const ivec&) {
return reduce_result_size;
};
auto w6 = mngr.spawn(prog, kn_reduce, cfg6,
opencl::in<int>{}, opencl::out<int>{result_size_6});
self->send(w6, move(arr6));
auto wg_size_as_int = static_cast<int>(max_wg_size);
ivec expected4{wg_size_as_int * 7, wg_size_as_int * 6, wg_size_as_int * 5,
wg_size_as_int * 4, wg_size_as_int * 3, wg_size_as_int * 2,
wg_size_as_int , 0};
self->receive(
[&](const ivec& result) {
check_vector_results("Passing size for the output", expected4, result);
}, others >> wrong_msg
); );
auto dev6 = mngr.get_device_if([](const device& d) {
return d.get_device_type() != cpu;
});
if (dev6) {
// test for manuel return size selection (max workgroup size 1d)
auto max_wg_size = min(dev6->get_max_work_item_sizes()[0], size_t{512});
auto reduce_buffer_size = static_cast<size_t>(max_wg_size) * 8;
auto reduce_local_size = static_cast<size_t>(max_wg_size);
auto reduce_work_groups = reduce_buffer_size / reduce_local_size;
auto reduce_global_size = reduce_buffer_size;
auto reduce_result_size = reduce_work_groups;
ivec arr6(reduce_buffer_size);
int n = static_cast<int>(arr6.capacity());
generate(arr6.begin(), arr6.end(), [&]{ return --n; });
opencl::spawn_config cfg6{dims{reduce_global_size},
dims{ }, // no offset
dims{reduce_local_size}};
auto get_result_size_6 = [reduce_result_size](const ivec&) {
return reduce_result_size;
};
auto w6 = mngr.spawn(mngr.create_program(kernel_source_reduce, "", *dev6),
kernel_name_reduce, cfg6,
opencl::in<ivec>{}, opencl::out<ivec>{get_result_size_6});
self->send(w6, move(arr6));
auto wg_size_as_int = static_cast<int>(max_wg_size);
ivec expected4{wg_size_as_int * 7, wg_size_as_int * 6, wg_size_as_int * 5,
wg_size_as_int * 4, wg_size_as_int * 3, wg_size_as_int * 2,
wg_size_as_int , 0};
self->receive(
[&](const ivec& result) {
check_vector_results("Passing size for the output", expected4, result);
},
others >> [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
}
);
}
// calculator function for getting the size of the output // calculator function for getting the size of the output
auto get_result_size_7 = [=](const ivec&) { auto result_size_7 = [](const ivec&) {
return problem_size; return problem_size;
}; };
// constant memory arguments // constant memory arguments
const ivec arr7{problem_size}; const ivec arr7{problem_size};
auto w7 = mngr.spawn(kernel_source_const, kernel_name_const, auto w7 = mngr.spawn(kernel_source, kn_const,
opencl::spawn_config{dims{problem_size}}, opencl::spawn_config{dims{problem_size}},
opencl::in<ivec>{}, opencl::in<int>{},
opencl::out<ivec>{get_result_size_7}); opencl::out<int>{result_size_7});
self->send(w7, move(arr7)); self->send(w7, move(arr7));
ivec expected5(problem_size); ivec expected5(problem_size);
fill(begin(expected5), end(expected5), static_cast<int>(problem_size)); fill(begin(expected5), end(expected5), static_cast<int>(problem_size));
self->receive( self->receive(
[&](const ivec& result) { [&](const ivec& result) {
check_vector_results("Using const input argument", expected5, result); check_vector_results("Using const input argument", expected5, result);
}, }, others >> wrong_msg
others >> [&](message_view& x) -> result<message> { );
CAF_ERROR("unexpected message" << x.content().stringify()); }
return sec::unexpected_message;
} void test_arguments(actor_system& sys) {
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device_if([](const device_ptr){ return true; });
CAF_REQUIRE(opt);
auto dev = *opt;
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
const ivec expected1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto w1 = mngr.spawn(mngr.create_program(kernel_source, "", dev), kn_matrix,
opencl::spawn_config{dims{matrix_size, matrix_size}},
opencl::in<int>{}, opencl::out<int>{});
self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
self->receive (
[&](const ivec& result) {
check_vector_results("arguments: from in to out", expected1, result);
}, others >> wrong_msg
); );
// test in_out argument type
ivec input9 = make_iota_vector<int>(problem_size); ivec input9 = make_iota_vector<int>(problem_size);
ivec expected9{input9}; ivec expected9{input9};
transform(begin(expected9), end(expected9), begin(expected9), for_each(begin(expected9), end(expected9), [](int& val){ val *= 2; });
[](const int& val){ return val * 2; }); auto w9 = mngr.spawn(kernel_source, kn_inout,
auto w9 = mngr.spawn(kernel_source_inout, kernel_name_inout,
spawn_config{dims{problem_size}}, spawn_config{dims{problem_size}},
opencl::in_out<ivec>{}); opencl::in_out<int>{});
self->send(w9, move(input9)); self->send(w9, move(input9));
self->receive( self->receive(
[&](const ivec& result) { [&](const ivec& result) {
check_vector_results("Testing in_out arugment", expected9, result); check_vector_results("Testing in_out arugment", expected9, result);
}, }, others >> wrong_msg
others >> [&](message_view& x) -> result<message> { );
CAF_ERROR("unexpected message" << x.content().stringify()); ivec input10 = make_iota_vector<int>(problem_size);
return sec::unexpected_message; ivec expected10{input10};
for_each(begin(expected10), end(expected10), [](int& val){ val *= 2; });
auto result_size_10 = [=](const ivec& input) { return input.size(); };
auto w10 = mngr.spawn(kernel_source, kn_scratch,
spawn_config{dims{problem_size}},
opencl::in_out<int>{},
opencl::scratch<int>{result_size_10});
self->send(w10, move(input10));
self->receive(
[&](const ivec& result) {
check_vector_results("Testing buffer arugment", expected10, result);
}, others >> wrong_msg
);
// test local
size_t la_global = 256;
size_t la_local = la_global / 2;
ivec input_local = make_iota_vector<int>(la_global);
ivec expected_local{input_local};
auto last = 0;
for (size_t i = 0; i < la_global; ++i) {
if (i == la_local) last = 0;
auto tmp = expected_local[i];
expected_local[i] = last;
last += tmp;
}
auto work_local = mngr.spawn(kernel_source, kn_local,
spawn_config{dims{la_global}, {}, dims{la_local}},
opencl::in_out<int>{},
opencl::local<int>{la_local});
self->send(work_local, std::move(input_local));
self->receive(
[&](const ivec& result) {
check_vector_results("Testing local arugment", expected_local, result);
}
);
// Same test, different argument order
input_local = make_iota_vector<int>(la_global);
work_local = mngr.spawn(kernel_source, kn_order,
spawn_config{dims{la_global}, {}, dims{la_local}},
opencl::local<int>{la_local},
opencl::in_out<int>{});
self->send(work_local, std::move(input_local));
self->receive(
[&](const ivec& result) {
check_vector_results("Testing local arugment", expected_local, result);
}
);
// Test private argument
ivec input_private = make_iota_vector<int>(problem_size);
int val_private = 42;
ivec expected_private{input_private};
for_each(begin(expected_private), end(expected_private),
[val_private](int& val){ val += val_private; });
auto worker_private = mngr.spawn(kernel_source, kn_private,
spawn_config{dims{problem_size}},
opencl::in_out<int>{},
opencl::priv<int>{val_private});
self->send(worker_private, std::move(input_private));
self->receive(
[&](const ivec& result) {
check_vector_results("Testing private arugment", expected_private,
result);
} }
); );
} }
CAF_TEST(test_opencl) { CAF_TEST(opencl_basics) {
actor_system_config cfg; actor_system_config cfg;
cfg.load<opencl::manager>() cfg.load<opencl::manager>()
.add_message_type<ivec>("int_vector") .add_message_type<ivec>("int_vector")
...@@ -418,3 +582,535 @@ CAF_TEST(test_opencl) { ...@@ -418,3 +582,535 @@ CAF_TEST(test_opencl) {
test_opencl(system); test_opencl(system);
system.await_all_actors_done(); system.await_all_actors_done();
} }
CAF_TEST(opencl_arguments) {
actor_system_config cfg;
cfg.load<opencl::manager>()
.add_message_type<ivec>("int_vector")
.add_message_type<matrix_type>("square_matrix");
actor_system system{cfg};
test_arguments(system);
system.await_all_actors_done();
}
CAF_TEST(opencl_mem_refs) {
actor_system_config cfg;
cfg.load<opencl::manager>();
actor_system system{cfg};
auto& mngr = system.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
// global arguments
vector<uint32_t> input{1, 2, 3, 4};
auto buf_1 = dev->global_argument(input, buffer_type::input_output);
CAF_CHECK_EQUAL(buf_1.size(), input.size());
auto res_1 = buf_1.data();
CAF_CHECK(res_1);
CAF_CHECK_EQUAL(res_1->size(), input.size());
check_vector_results("Testing mem_ref", input, *res_1);
auto res_2 = buf_1.data(2ul);
CAF_CHECK(res_2);
CAF_CHECK_EQUAL(res_2->size(), 2ul);
CAF_CHECK_EQUAL((*res_2)[0], input[0]);
CAF_CHECK_EQUAL((*res_2)[1], input[1]);
vector<uint32_t> new_input{1,2,3,4,5};
buf_1 = dev->global_argument(new_input, buffer_type::input_output);
CAF_CHECK_EQUAL(buf_1.size(), new_input.size());
auto res_3 = buf_1.data();
CAF_CHECK(res_3);
mem_ref<uint32_t> buf_2{std::move(buf_1)};
CAF_CHECK_EQUAL(buf_2.size(), new_input.size());
auto res_4 = buf_2.data();
CAF_CHECK(res_4);
buf_2.reset();
auto res_5 = buf_2.data();
CAF_CHECK(!res_5);
}
CAF_TEST(opencl_argument_info) {
using base_t = int;
using in_arg_t = caf::detail::type_list<opencl::in<base_t>>;
using in_arg_info_t = typename cl_arg_info_list<in_arg_t>::type;
using in_arg_wrap_t = typename caf::detail::tl_head<in_arg_info_t>::type;
static_assert(in_arg_wrap_t::in_pos == 0, "In-index for `in` wrong.");
static_assert(in_arg_wrap_t::out_pos == -1, "Out-index for `in` wrong.");
using out_arg_t = caf::detail::type_list<opencl::out<base_t>>;
using out_arg_info_t = typename cl_arg_info_list<out_arg_t>::type;
using out_arg_wrap_t = typename caf::detail::tl_head<out_arg_info_t>::type;
static_assert(out_arg_wrap_t::in_pos == -1, "In-index for `out` wrong.");
static_assert(out_arg_wrap_t::out_pos == 0, "Out-index for `out` wrong.");
using io_arg_t = caf::detail::type_list<opencl::in_out<base_t>>;
using io_arg_info_t = typename cl_arg_info_list<io_arg_t>::type;
using io_arg_wrap_t = typename caf::detail::tl_head<io_arg_info_t>::type;
static_assert(io_arg_wrap_t::in_pos == 0, "In-index for `in_out` wrong.");
static_assert(io_arg_wrap_t::out_pos == 0, "Out-index for `in_out` wrong.");
using arg_list_t = caf::detail::type_list<opencl::in<base_t>,
opencl::out<base_t>,
opencl::local<base_t>,
opencl::in_out<base_t>,
opencl::priv<base_t>,
opencl::priv<base_t, val>>;
using arg_info_list_t = typename cl_arg_info_list<arg_list_t>::type;
using arg_info_0_t = typename caf::detail::tl_at<arg_info_list_t,0>::type;
static_assert(arg_info_0_t::in_pos == 0, "In-index for `in` wrong.");
static_assert(arg_info_0_t::out_pos == -1, "Out-index for `in` wrong.");
using arg_info_1_t = typename caf::detail::tl_at<arg_info_list_t,1>::type;
static_assert(arg_info_1_t::in_pos == -1, "In-index for `out` wrong.");
static_assert(arg_info_1_t::out_pos == 0, "Out-index for `out` wrong.");
using arg_info_2_t = typename caf::detail::tl_at<arg_info_list_t,2>::type;
static_assert(arg_info_2_t::in_pos == -1, "In-index for `local` wrong.");
static_assert(arg_info_2_t::out_pos == -1, "Out-index for `local` wrong.");
using arg_info_3_t = typename caf::detail::tl_at<arg_info_list_t,3>::type;
static_assert(arg_info_3_t::in_pos == 1, "In-index for `in_out` wrong.");
static_assert(arg_info_3_t::out_pos == 1, "Out-index for `in_out` wrong.");
using arg_info_4_t = typename caf::detail::tl_at<arg_info_list_t,4>::type;
static_assert(arg_info_4_t::in_pos == -1, "In-index for `priv` wrong.");
static_assert(arg_info_4_t::out_pos == -1, "Out-index for `priv` wrong.");
using arg_info_5_t = typename caf::detail::tl_at<arg_info_list_t,5>::type;
static_assert(arg_info_5_t::in_pos == 2, "In-index for `priv` wrong.");
static_assert(arg_info_5_t::out_pos == -1, "Out-index for `priv` wrong.");
// gives the test some output.
CAF_CHECK_EQUAL(true, true);
}
void test_in_val_out_val(actor_system& sys) {
CAF_MESSAGE("Testing in: val -> out: val ");
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, conf, in<int>{}, out<int>{});
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)", res1, result);
}, others >> wrong_msg);
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, conf,
in<int>{}, out<int>{});
self->send(w2, make_iota_vector<int>(matrix_size * matrix_size));
self->receive([&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors"
" (kernel passed directly)", res1, result);
}, others >> wrong_msg);
// Wrap message in user-defined type and use mapping functions
const matrix_type res2(move(res1));
auto map_arg = [](message& msg) -> optional<message> {
return msg.apply([](matrix_type& mx) {
return make_message(move(mx.data()));
});
};
auto map_res = [](ivec result) -> message {
return make_message(matrix_type{move(result)});
};
auto w3 = mngr.spawn(prog, kn_matrix, conf, map_arg, map_res,
in<int, val>{}, out<int, val>{});
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)",
res2.data(), result.data());
}, others >> wrong_msg);
// create program with opencl compiler flags
auto prog2 = mngr.create_program(kernel_source_compiler_flag, compiler_flag);
spawn_config conf2{dims{array_size}};
auto w4 = mngr.spawn(prog2, kn_compiler_flag, conf2,
in<int>{}, out<int>{});
self->send(w4, make_iota_vector<int>(array_size));
auto res3 = make_iota_vector<int>(array_size);
self->receive([&](const ivec& result) {
check_vector_results("Passing compiler flags", res3, result);
}, others >> wrong_msg);
// test for manuel return size selection (max workgroup size 1d)
auto max_wg_size = min(dev->get_max_work_item_sizes()[0], size_t{512});
auto reduce_buffer_size = static_cast<size_t>(max_wg_size) * 8;
auto reduce_local_size = static_cast<size_t>(max_wg_size);
auto reduce_work_groups = reduce_buffer_size / reduce_local_size;
auto reduce_global_size = reduce_buffer_size;
auto reduce_result_size = reduce_work_groups;
ivec input(reduce_buffer_size);
int n = static_cast<int>(input.capacity());
generate(input.begin(), input.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
auto res_size = [&](const ivec&) { return reduce_result_size; };
auto w5 = mngr.spawn(prog, kn_reduce, conf3,
in<int>{}, out<int>{res_size});
self->send(w5, move(input));
auto wg_size_as_int = static_cast<int>(max_wg_size);
ivec res4{
wg_size_as_int * 7, wg_size_as_int * 6, wg_size_as_int * 5,
wg_size_as_int * 4, wg_size_as_int * 3, wg_size_as_int * 2,
wg_size_as_int , 0
};
self->receive([&](const ivec& result) {
check_vector_results("Passing size for the output", res4, result);
}, others >> wrong_msg);
// calculator function for getting the size of the output
auto res_size2 = [](const ivec&) { return problem_size; };
// constant memory arguments
const ivec input2{problem_size};
auto w6 = mngr.spawn(kernel_source, kn_const,
spawn_config{dims{problem_size}},
in<int>{}, out<int>{res_size2});
self->send(w6, move(input2));
ivec res5(problem_size);
fill(begin(res5), end(res5), static_cast<int>(problem_size));
self->receive([&](const ivec& result) {
check_vector_results("Using const input argument", res5, result);
}, others >> wrong_msg);
}
void test_in_val_out_mref(actor_system& sys) {
CAF_MESSAGE("Testing in: val -> out: mref");
// setup
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
// tests
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, conf, in<int>{}, out<int, mref>{});
self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
self->receive([&](iref& result) {
check_mref_results("Simple matrix multiplication using vectors"
" (kernel wrapped in program)", res1, result);
}, others >> wrong_msg);
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, conf,
in<int>{}, out<int, mref>{});
self->send(w2, make_iota_vector<int>(matrix_size * matrix_size));
self->receive([&](iref& result) {
check_mref_results("Simple matrix multiplication using vectors"
" (kernel passed directly)", res1, result);
}, others >> wrong_msg);
// test for manuel return size selection (max workgroup size 1d)
auto max_wg_size = min(dev->get_max_work_item_sizes()[0], size_t{512});
auto reduce_buffer_size = static_cast<size_t>(max_wg_size) * 8;
auto reduce_local_size = static_cast<size_t>(max_wg_size);
auto reduce_work_groups = reduce_buffer_size / reduce_local_size;
auto reduce_global_size = reduce_buffer_size;
auto reduce_result_size = reduce_work_groups;
ivec input(reduce_buffer_size);
int n = static_cast<int>(input.capacity());
generate(input.begin(), input.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
auto res_size = [&](const ivec&) { return reduce_result_size; };
auto w5 = mngr.spawn(prog, kn_reduce, conf3,
in<int>{}, out<int, mref>{res_size});
self->send(w5, move(input));
auto wg_size_as_int = static_cast<int>(max_wg_size);
ivec res4{wg_size_as_int * 7, wg_size_as_int * 6, wg_size_as_int * 5,
wg_size_as_int * 4, wg_size_as_int * 3, wg_size_as_int * 2,
wg_size_as_int * 1, wg_size_as_int * 0};
self->receive([&](iref& result) {
check_mref_results("Passing size for the output", res4, result);
}, others >> wrong_msg);
}
void test_in_mref_out_val(actor_system& sys) {
CAF_MESSAGE("Testing in: mref -> out: val ");
// setup
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
// tests
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, conf, in<int, mref>{}, out<int>{});
auto matrix1 = make_iota_vector<int>(matrix_size * matrix_size);
auto input1 = dev->global_argument(matrix1);
self->send(w1, input1);
self->receive([&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors"
" (kernel wrapped in program)", res1, result);
}, others >> wrong_msg);
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, conf,
in<int, mref>{}, out<int, val>{});
self->send(w2, input1);
self->receive([&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors"
" (kernel passed directly)", res1, result);
}, others >> wrong_msg);
// test for manuel return size selection (max workgroup size 1d)
auto max_wg_size = min(dev->get_max_work_item_sizes()[0], size_t{512});
auto reduce_buffer_size = static_cast<size_t>(max_wg_size) * 8;
auto reduce_local_size = static_cast<size_t>(max_wg_size);
auto reduce_work_groups = reduce_buffer_size / reduce_local_size;
auto reduce_global_size = reduce_buffer_size;
auto reduce_result_size = reduce_work_groups;
ivec values(reduce_buffer_size);
int n = static_cast<int>(values.capacity());
generate(values.begin(), values.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
auto res_size = [&](const iref&) { return reduce_result_size; };
auto w5 = mngr.spawn(prog, kn_reduce, conf3,
in<int, mref>{}, out<int>{res_size});
auto input2 = dev->global_argument(values);
self->send(w5, input2);
auto multiplier = static_cast<int>(max_wg_size);
ivec res4{multiplier * 7, multiplier * 6, multiplier * 5,
multiplier * 4, multiplier * 3, multiplier * 2,
multiplier * 1, multiplier * 0};
self->receive([&](const ivec& result) {
check_vector_results("Passing size for the output", res4, result);
}, others >> wrong_msg);
}
void test_in_mref_out_mref(actor_system& sys) {
CAF_MESSAGE("Testing in: mref -> out: mref");
// setup
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
// tests
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, conf,
in<int, mref>{}, out<int, mref>{});
auto matrix1 = make_iota_vector<int>(matrix_size * matrix_size);
auto input1 = dev->global_argument(matrix1);
self->send(w1, input1);
self->receive([&](iref& result) {
check_mref_results("Simple matrix multiplication using vectors"
" (kernel wrapped in program)", res1, result);
}, others >> wrong_msg);
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, conf,
in<int, mref>{}, out<int, mref>{});
self->send(w2, input1);
self->receive([&](iref& result) {
check_mref_results("Simple matrix multiplication using vectors"
" (kernel passed directly)", res1, result);
}, others >> wrong_msg);
// test for manuel return size selection (max workgroup size 1d)
auto max_wg_size = min(dev->get_max_work_item_sizes()[0], size_t{512});
auto reduce_buffer_size = static_cast<size_t>(max_wg_size) * 8;
auto reduce_local_size = static_cast<size_t>(max_wg_size);
auto reduce_work_groups = reduce_buffer_size / reduce_local_size;
auto reduce_global_size = reduce_buffer_size;
auto reduce_result_size = reduce_work_groups;
ivec values(reduce_buffer_size);
int n = static_cast<int>(values.capacity());
generate(values.begin(), values.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
auto res_size = [&](const iref&) { return reduce_result_size; };
auto w5 = mngr.spawn(prog, kn_reduce, conf3,
in<int, mref>{}, out<int, mref>{res_size});
auto input2 = dev->global_argument(values);
self->send(w5, input2);
auto multiplier = static_cast<int>(max_wg_size);
ivec res4{multiplier * 7, multiplier * 6, multiplier * 5,
multiplier * 4, multiplier * 3, multiplier * 2,
multiplier , 0};
self->receive([&](iref& result) {
check_mref_results("Passing size for the output", res4, result);
}, others >> wrong_msg);
}
void test_varying_arguments(actor_system& sys) {
CAF_MESSAGE("Testing varying argument order "
"(Might fail on some integrated GPUs)");
// setup
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
// tests
size_t size = 23;
spawn_config conf{dims{size}};
auto input1 = make_iota_vector<int>(size);
auto input2 = dev->global_argument(input1);
auto w1 = mngr.spawn(prog, kn_varying, conf,
in<int>{}, out<int>{}, in<int>{}, out<int>{});
self->send(w1, input1, input1);
self->receive([&](const ivec& res1, const ivec& res2) {
check_vector_results("Varying args (vec only), output 1", input1, res1);
check_vector_results("Varying args (vec only), output 2", input1, res2);
}, others >> wrong_msg);
auto w2 = mngr.spawn(prog, kn_varying, conf,
in<int,mref>{}, out<int>{},
in<int>{}, out<int,mref>{});
self->send(w2, input2, input1);
self->receive([&](const ivec& res1, iref& res2) {
check_vector_results("Varying args (vec), output 1", input1, res1);
check_mref_results("Varying args (ref), output 2", input1, res2);
}, others >> wrong_msg);
}
void test_inout(actor_system& sys) {
CAF_MESSAGE("Testing in_out arguments");
// setup
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
// tests
ivec input = make_iota_vector<int>(problem_size);
auto input2 = dev->global_argument(input);
auto input3 = dev->global_argument(input);
ivec res{input};
for_each(begin(res), end(res), [](int& val){ val *= 2; });
auto conf = spawn_config{dims{problem_size}};
auto w1 = mngr.spawn(kernel_source, kn_inout, conf,
in_out<int,val,val>{});
self->send(w1, input);
self->receive([&](const ivec& result) {
check_vector_results("Testing in_out (val -> val)", res, result);
}, others >> wrong_msg);
auto w2 = mngr.spawn(kernel_source, kn_inout, conf,
in_out<int,val,mref>{});
self->send(w2, input);
self->receive([&](iref& result) {
check_mref_results("Testing in_out (val -> mref)", res, result);
}, others >> wrong_msg);
auto w3 = mngr.spawn(kernel_source, kn_inout, conf,
in_out<int,mref,val>{});
self->send(w3, input2);
self->receive([&](const ivec& result) {
check_vector_results("Testing in_out (mref -> val)", res, result);
}, others >> wrong_msg);
auto w4 = mngr.spawn(kernel_source, kn_inout, conf,
in_out<int,mref,mref>{});
self->send(w4, input3);
self->receive([&](iref& result) {
check_mref_results("Testing in_out (mref -> mref)", res, result);
}, others >> wrong_msg);
}
void test_priv(actor_system& sys) {
CAF_MESSAGE("Testing priv argument");
// setup
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
// tests
spawn_config conf{dims{problem_size}};
ivec input = make_iota_vector<int>(problem_size);
int value = 42;
ivec res{input};
for_each(begin(res), end(res), [&](int& val){ val += value; });
auto w1 = mngr.spawn(kernel_source, kn_private, conf,
in_out<int>{}, priv<int>{value});
self->send(w1, input);
self->receive([&](const ivec& result) {
check_vector_results("Testing hidden private arugment", res, result);
}, others >> wrong_msg);
auto w2 = mngr.spawn(kernel_source, kn_private, conf,
in_out<int>{}, priv<int,val>{});
self->send(w2, input, value);
self->receive([&](const ivec& result) {
check_vector_results("Testing val private arugment", res, result);
}, others >> wrong_msg);
}
void test_local(actor_system& sys) {
CAF_MESSAGE("Testing local argument");
// setup
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device(0);
CAF_REQUIRE(opt);
auto dev = *opt;
auto prog = mngr.create_program(kernel_source, "", dev);
scoped_actor self{sys};
auto wrong_msg = [&](message_view& x) -> result<message> {
CAF_ERROR("unexpected message" << x.content().stringify());
return sec::unexpected_message;
};
// tests
size_t global_size = 256;
size_t local_size = global_size / 2;
ivec res = make_iota_vector<int>(global_size);
auto last = 0;
for (size_t i = 0; i < global_size; ++i) {
if (i == local_size) last = 0;
auto tmp = res[i];
res[i] = last;
last += tmp;
}
auto conf = spawn_config{dims{global_size}, {}, dims{local_size}};
auto w = mngr.spawn(kernel_source, kn_local, conf,
in_out<int>{}, local<int>{local_size});
self->send(w, make_iota_vector<int>(global_size));
self->receive([&](const ivec& result) {
check_vector_results("Testing local arugment", res, result);
}, others >> wrong_msg);
// Same test, different argument order
w = mngr.spawn(kernel_source, kn_order, conf,
local<int>{local_size}, in_out<int>{});
self->send(w, make_iota_vector<int>(global_size));
self->receive([&](const ivec& result) {
check_vector_results("Testing local arugment", res, result);
}, others >> wrong_msg);
}
CAF_TEST(opencl_opencl_actor) {
actor_system_config cfg;
cfg.load<opencl::manager>()
.add_message_type<ivec>("int_vector")
.add_message_type<matrix_type>("square_matrix");
actor_system system{cfg};
test_in_val_out_val(system);
test_in_val_out_mref(system);
test_in_mref_out_val(system);
test_in_mref_out_mref(system);
test_varying_arguments(system);
test_inout(system);
test_priv(system);
test_local(system);
system.await_all_actors_done();
}
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