Commit 36b3f4f3 authored by Joseph Noir's avatar Joseph Noir

Adapt to actor_system update

parent 785153d3
......@@ -4,13 +4,14 @@ project(caf_opencl C CXX)
# get header files; only needed by CMake generators,
# e.g., for creating proper Xcode projects
file(GLOB LIBCAF_OPENCL_HDRS "caf/opencl/*.hpp")
file(GLOB LIBCAF_OPENCL_HDRS "caf/opencl/detail/*.hpp")
add_custom_target(libcaf_opencl)
# list cpp files excluding platform-dependent files
set (LIBCAF_OPENCL_SRCS
src/global.cpp
src/metainfo.cpp
src/manager.cpp
src/program.cpp
src/opencl_err.cpp
src/platform.cpp
......@@ -44,4 +45,3 @@ link_directories(${LD_DIRS})
include_directories(. ${INCLUDE_DIRS})
# install includes
install(DIRECTORY caf/ DESTINATION include/caf FILES_MATCHING PATTERN "*.hpp")
install(DIRECTORY cppa/ DESTINATION include/cppa FILES_MATCHING PATTERN "*.hpp")
......@@ -27,7 +27,6 @@
#include "caf/all.hpp"
#include "caf/channel.hpp"
#include "caf/intrusive_ptr.hpp"
#include "caf/detail/int_list.hpp"
......@@ -45,7 +44,7 @@
namespace caf {
namespace opencl {
class metainfo;
class manager;
template <class List>
struct function_sig_from_outputs;
......@@ -64,7 +63,7 @@ struct command_sig_from_outputs<T, detail::type_list<Ts...>> {
};
template <class... Ts>
class actor_facade : public abstract_actor {
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;
......@@ -73,7 +72,7 @@ public:
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<maybe<message> (message&)>;
using input_mapping = std::function<optional<message> (message&)>;
using output_wrapped_types =
typename detail::tl_filter<arg_types, is_output_arg>::type;
......@@ -90,40 +89,57 @@ public:
using command_type =
typename command_sig_from_outputs<actor_facade, output_types>::type;
static intrusive_ptr<actor_facade> create(const program& prog,
const char* kernel_name,
const spawn_config& config,
input_mapping map_args,
output_mapping map_result,
Ts&&... xs) {
if (config.dimensions().empty()) {
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_LOGF_ERROR(str);
CAF_LOG_ERROR(str);
throw std::runtime_error(str);
}
auto check_vec = [&](const dim_vec& vec, const char* name) {
if (! vec.empty() && vec.size() != config.dimensions().size()) {
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_LOGF_ERROR(oss.str());
CAF_LOG_ERROR(CAF_ARG(oss.str()));
throw std::runtime_error(oss.str());
}
};
check_vec(config.offsets(), "offsets");
check_vec(config.local_dimensions(), "local dimensions");
cl_int err = 0;
kernel_ptr kernel;
kernel.reset(clCreateKernel(prog.program_.get(), kernel_name, &err), false);
if (err != CL_SUCCESS)
return nullptr;
return new actor_facade(prog, kernel, config,
std::move(map_args), std::move(map_result),
std::forward_as_tuple(xs...));
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...));
} else {
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(const actor_addr &sender, message_id mid, message content,
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);
......@@ -135,14 +151,15 @@ public:
if (! content.match_elements(input_types{})) {
return;
}
response_promise hdl{this->address(), sender, mid.response_id()};
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>(hdl, this,
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),
......@@ -151,21 +168,29 @@ public:
cmd->enqueue();
}
actor_facade(const program& prog, kernel_ptr kernel,
const spawn_config& config,
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->msg, eu);
}
actor_facade(actor_config actor_cfg,
const program& prog, kernel_ptr kernel,
const spawn_config& spawn_cfg,
input_mapping map_args, output_mapping map_result,
std::tuple<Ts...> xs)
: kernel_(kernel),
: monitorable_actor(actor_cfg),
kernel_(kernel),
program_(prog.program_),
context_(prog.context_),
queue_(prog.queue_),
config_(config),
spawn_cfg_(spawn_cfg),
map_args_(std::move(map_args)),
map_results_(std::move(map_result)),
argument_types_(xs) {
CAF_LOG_TRACE("id: " << this->id());
default_output_size_ = std::accumulate(config_.dimensions().begin(),
config_.dimensions().end(),
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>{});
}
......@@ -258,7 +283,7 @@ public:
program_ptr program_;
context_ptr context_;
command_queue_ptr queue_;
spawn_config config_;
spawn_config spawn_cfg_;
input_mapping map_args_;
output_mapping map_results_;
std::tuple<Ts...> argument_types_;
......
......@@ -21,6 +21,6 @@
#ifndef CAF_OPENCL_ALL_HPP
#define CAF_OPENCL_ALL_HPP
#include "caf/opencl/spawn_cl.hpp"
#include "caf/opencl/manager.hpp"
#endif // CAF_OPENCL_ALL_HPP
......@@ -25,7 +25,7 @@
#include <type_traits>
#include "caf/message.hpp"
#include "caf/maybe.hpp"
#include "caf/optional.hpp"
namespace caf {
namespace opencl {
......@@ -56,7 +56,7 @@ struct out {
out() { }
template <class F>
out(F fun) {
fun_ = [fun](message& msg) -> maybe<size_t> {
fun_ = [fun](message& msg) -> optional<size_t> {
auto res = msg.apply(fun);
size_t result;
if (res) {
......@@ -66,10 +66,10 @@ struct out {
return none;
};
}
maybe<size_t> operator()(message& msg) const {
return fun_ ? fun_(msg) : 0;
optional<size_t> operator()(message& msg) const {
return fun_ ? fun_(msg) : 0UL;
}
std::function<maybe<size_t> (message&)> fun_;
std::function<optional<size_t> (message&)> fun_;
};
......
......@@ -26,10 +26,11 @@
#include <algorithm>
#include <functional>
#include "caf/logger.hpp"
#include "caf/actor_cast.hpp"
#include "caf/abstract_actor.hpp"
#include "caf/response_promise.hpp"
#include "caf/detail/logging.hpp"
#include "caf/detail/scope_guard.hpp"
#include "caf/opencl/global.hpp"
......@@ -43,14 +44,14 @@ namespace opencl {
template <class FacadeType, class... Ts>
class command : public ref_counted {
public:
command(response_promise handle, intrusive_ptr<FacadeType> actor_facade,
command(std::tuple<strong_actor_ptr,message_id> handle,
strong_actor_ptr actor_facade,
std::vector<cl_event> events, std::vector<mem_ptr> input_buffers,
std::vector<mem_ptr> output_buffers, std::vector<size_t> result_sizes,
message msg)
: result_sizes_(result_sizes),
handle_(handle),
actor_facade_(actor_facade),
queue_(actor_facade->queue_),
mem_in_events_(std::move(events)),
input_buffers_(std::move(input_buffers)),
output_buffers_(std::move(output_buffers)),
......@@ -76,18 +77,21 @@ public:
auto data_or_nullptr = [](const dim_vec& vec) {
return vec.empty() ? nullptr : vec.data();
};
auto actor_facade =
static_cast<FacadeType*>(actor_cast<abstract_actor*>(actor_facade_));
// OpenCL expects cl_uint (unsigned int), hence the cast
cl_int err = clEnqueueNDRangeKernel(
queue_.get(), actor_facade_->kernel_.get(),
static_cast<cl_uint>(actor_facade_->config_.dimensions().size()),
data_or_nullptr(actor_facade_->config_.offsets()),
data_or_nullptr(actor_facade_->config_.dimensions()),
data_or_nullptr(actor_facade_->config_.local_dimensions()),
actor_facade->queue_.get(), actor_facade->kernel_.get(),
static_cast<cl_uint>(actor_facade->spawn_cfg_.dimensions().size()),
data_or_nullptr(actor_facade->spawn_cfg_.offsets()),
data_or_nullptr(actor_facade->spawn_cfg_.dimensions()),
data_or_nullptr(actor_facade->spawn_cfg_.local_dimensions()),
static_cast<cl_uint>(mem_in_events_.size()),
(mem_in_events_.empty() ? nullptr : mem_in_events_.data()), &event_k
);
if (err != CL_SUCCESS) {
CAF_LOGMF(CAF_ERROR, "clEnqueueNDRangeKernel: " << get_opencl_error(err));
CAF_LOG_ERROR("clEnqueueNDRangeKernel: "
<< CAF_ARG(get_opencl_error(err)));
clReleaseEvent(event_k);
this->deref();
return;
......@@ -96,15 +100,15 @@ public:
cl_event marker;
#if defined(__APPLE__)
err = clEnqueueMarkerWithWaitList(
queue_.get(),
actor_facade->queue_.get(),
static_cast<cl_uint>(mem_out_events_.size()),
mem_out_events_.data(), &marker
);
#else
err = clEnqueueMarker(queue_.get(), &marker);
err = clEnqueueMarker(actor_facade->queue_.get(), &marker);
#endif
if (err != CL_SUCCESS) {
CAF_LOGMF(CAF_ERROR, "clSetEventCallback: " << get_opencl_error(err));
CAF_LOG_ERROR("clSetEventCallback: " << CAF_ARG(get_opencl_error(err)));
clReleaseEvent(marker);
clReleaseEvent(event_k);
this->deref(); // callback is not set
......@@ -118,15 +122,15 @@ public:
},
this);
if (err != CL_SUCCESS) {
CAF_LOGMF(CAF_ERROR, "clSetEventCallback: " << get_opencl_error(err));
CAF_LOG_ERROR("clSetEventCallback: " << CAF_ARG(get_opencl_error(err)));
clReleaseEvent(marker);
clReleaseEvent(event_k);
this->deref(); // callback is not set
return;
}
err = clFlush(queue_.get());
err = clFlush(actor_facade->queue_.get());
if (err != CL_SUCCESS) {
CAF_LOGMF(CAF_ERROR, "clFlush: " << get_opencl_error(err));
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));
......@@ -135,9 +139,8 @@ public:
private:
std::vector<size_t> result_sizes_;
response_promise handle_;
intrusive_ptr<FacadeType> actor_facade_;
command_queue_ptr queue_;
std::tuple<strong_actor_ptr,message_id> handle_;
strong_actor_ptr actor_facade_;
std::vector<cl_event> mem_in_events_;
std::vector<cl_event> mem_out_events_;
std::vector<mem_ptr> input_buffers_;
......@@ -151,6 +154,8 @@ private:
template <long I, long... Is>
void enqueue_read_buffers(cl_event& kernel_done, detail::int_list<I, Is...>) {
auto actor_facade =
static_cast<FacadeType*>(actor_cast<abstract_actor*>(actor_facade_));
using container_type =
typename std::tuple_element<I, std::tuple<Ts...>>::type;
using value_type = typename container_type::value_type;
......@@ -158,7 +163,8 @@ private:
auto size = result_sizes_[I];
auto buffer_size = sizeof(value_type) * result_sizes_[I];
std::get<I>(result_buffers_).resize(size);
auto err = clEnqueueReadBuffer(queue_.get(), output_buffers_[I].get(),
auto err = clEnqueueReadBuffer(actor_facade->queue_.get(),
output_buffers_[I].get(),
CL_FALSE, 0, buffer_size,
std::get<I>(result_buffers_).data(),
1, &kernel_done, &event);
......@@ -172,12 +178,15 @@ private:
}
void handle_results() {
auto& map_fun = actor_facade_->map_results_;
auto actor_facade =
static_cast<FacadeType*>(actor_cast<abstract_actor*>(actor_facade_));
auto& map_fun = actor_facade->map_results_;
auto msg = map_fun ? apply_args(map_fun,
detail::get_indices(result_buffers_),
result_buffers_)
: message_from_results{}(result_buffers_);
handle_.deliver(std::move(msg));
get<0>(handle_)->enqueue(actor_facade_, get<1>(handle_), std::move(msg),
nullptr);
}
};
......
......@@ -17,68 +17,45 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_METAINFO_HPP
#define CAF_METAINFO_HPP
#ifndef CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
#define CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
#include <atomic>
#include <vector>
#include <algorithm>
#include <functional>
#include "caf/all.hpp"
#include "caf/config.hpp"
#include "caf/maybe.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/actor_facade.hpp"
namespace caf {
namespace opencl {
class metainfo : public detail::abstract_singleton {
friend class program;
friend class detail::singletons;
friend command_queue_ptr get_command_queue(uint32_t id);
public:
/// Get a list of all available devices. This is depricated, use the more specific
/// get_deivce and get_deivce_if functions.
/// (Returns only devices of the first discovered platform).
const std::vector<device>& get_devices() const CAF_DEPRECATED;
/// Get the device with id. These ids are assigned sequientally to all available devices.
const maybe<const device&> get_device(size_t id = 0) const;
/// Get the first device that satisfies the predicate.
/// The predicate should accept a `const device&` and return a bool;
template <class UnaryPredicate>
const maybe<const device&> get_device_if(UnaryPredicate p) const {
for (auto& pl : platforms_) {
for (auto& dev : pl.get_devices()) {
if (p(dev))
return dev;
}
}
return none;
namespace detail {
struct tuple_construct { };
template <class... Ts>
struct cl_spawn_helper {
using impl = opencl::actor_facade<Ts...>;
using map_in_fun = std::function<optional<message> (message&)>;
using map_out_fun = typename impl::output_mapping;
actor operator()(actor_config actor_cfg, const opencl::program& p,
const char* fn, const opencl::spawn_config& spawn_cfg,
Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg),
p, fn, spawn_cfg,
map_in_fun{}, map_out_fun{},
std::move(xs)...));
}
actor operator()(actor_config actor_cfg, const opencl::program& p,
const char* fn, const opencl::spawn_config& spawn_cfg,
map_in_fun map_input, map_out_fun map_output,
Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg),
p, fn, spawn_cfg,
std::move(map_input),
std::move(map_output),
std::move(xs)...));
}
/// Get metainfo instance.
static metainfo* instance();
private:
metainfo() = default;
void stop() override;
void initialize() override;
void dispose() override;
std::vector<platform> platforms_;
};
} // namespace detail
} // namespace opencl
} // namespace caf
#endif // CAF_METAINFO_HPP
#endif // CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
......@@ -29,11 +29,13 @@ namespace caf {
namespace opencl {
class program;
class manager;
class device {
public:
friend class program;
friend class manager;
public:
/// Intialize a new device in a context using a sepcific device_id
static device create(context_ptr context, device_ptr device_id, unsigned id);
/// Get the id assigned by caf
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_MANAGER_HPP
#define CAF_OPENCL_MANAGER_HPP
#include <atomic>
#include <vector>
#include <algorithm>
#include <functional>
#include "caf/optional.hpp"
#include "caf/config.hpp"
#include "caf/actor_system.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/actor_facade.hpp"
#include "caf/opencl/detail/spawn_helper.hpp"
namespace caf {
namespace opencl {
class manager : public actor_system::module {
public:
friend class program;
friend class actor_system;
friend command_queue_ptr get_command_queue(uint32_t id);
manager(const manager&) = delete;
manager& operator=(const manager&) = delete;
/// Get the device with id, which is assigned sequientally.
const optional<const device&> get_device(size_t id = 0) const;
/// Get the first device that satisfies the predicate.
/// The predicate should accept a `const device&` and return a bool;
template <class UnaryPredicate>
const optional<const device&> get_device_if(UnaryPredicate p) const {
for (auto& pl : platforms_) {
for (auto& dev : pl.get_devices()) {
if (p(dev))
return dev;
}
}
return none;
}
void start() override;
void stop() override;
void init(actor_system_config&) override;
id_t id() const override;
void* subtype_ptr() override;
static actor_system::module* make(actor_system& sys,
caf::detail::type_list<>);
// OpenCL functionality
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
program create_program(const char* kernel_source,
const char* options = nullptr, uint32_t device_id = 0);
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
program create_program(const char* kernel_source,
const char* options, const device& dev);
/// 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& prog,
const char* fname,
const opencl::spawn_config& config,
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, config,
std::move(x), std::move(xs)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template <class T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn(const char* source,
const char* fname,
const opencl::spawn_config& config,
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, config,
std::move(x), std::move(xs)...);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template <class Fun, class... Ts>
actor spawn(const opencl::program& prog,
const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args,
Fun map_result,
Ts... xs) {
detail::cl_spawn_helper<Ts...> f;
return f(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>
actor spawn(const char* source,
const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (message&)> map_args,
Fun map_result,
Ts... xs) {
detail::cl_spawn_helper<Ts...> f;
return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, config,
std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
protected:
manager(actor_system& sys);
~manager();
private:
actor_system& system_;
std::vector<platform> platforms_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_MANAGER_HPP
......@@ -24,7 +24,7 @@
#include "caf/opencl/global.hpp"
#include "caf/detail/logging.hpp"
#include "caf/logger.hpp"
#define CAF_CLF(funname) #funname , funname
......
......@@ -26,10 +26,9 @@ namespace caf {
namespace opencl {
class platform {
public:
friend class program;
public:
inline const std::vector<device>& get_devices() const;
inline const std::string& get_name() const;
inline const std::string& get_vendor() const;
......
......@@ -35,29 +35,19 @@ class actor_facade;
/// @brief A wrapper for OpenCL's cl_program.
class program {
public:
friend class manager;
template <class... Ts>
friend class actor_facade;
public:
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
static program create(const char* kernel_source,
const char* options = nullptr, uint32_t device_id = 0);
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
static program create(const char* kernel_source,
const char* options, const device& dev);
private:
program(context_ptr context, command_queue_ptr queue, program_ptr prog);
program(context_ptr context, command_queue_ptr queue, program_ptr prog,
std::map<std::string,kernel_ptr> available_kernels);
context_ptr context_;
program_ptr program_;
command_queue_ptr queue_;
// save CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
std::map<std::string,kernel_ptr> available_kernels_;
};
} // namespace opencl
......
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_SPAWN_CL_HPP
#define CAF_SPAWN_CL_HPP
#include <algorithm>
#include <functional>
#include "caf/maybe.hpp"
#include "caf/actor_cast.hpp"
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/metainfo.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/actor_facade.hpp"
#include "caf/opencl/spawn_config.hpp"
namespace caf {
namespace detail {
struct tuple_construct { };
template <class... Ts>
struct cl_spawn_helper {
using impl = opencl::actor_facade<Ts...>;
using map_in_fun = std::function<maybe<message> (message&)>;
using map_out_fun = typename impl::output_mapping;
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg, Ts&&... xs) const {
return actor_cast<actor>(impl::create(p, fn, cfg,
map_in_fun{}, map_out_fun{},
std::move(xs)...));
}
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg,
map_in_fun map_input, map_out_fun map_output,
Ts&&... xs) const {
return actor_cast<actor>(impl::create(p, fn, cfg, std::move(map_input),
std::move(map_output),
std::move(xs)...));
}
// used by the deprecated spawn_helper
template <class Tuple, long... Is>
actor operator()(tuple_construct,
const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg,
Tuple&& xs,
detail::int_list<Is...>) const {
return actor_cast<actor>(impl::create(p, fn, cfg,
map_in_fun{}, map_out_fun{},
std::move(std::get<Is>(xs))...));
}
template <class Tuple, long... Is>
actor operator()(tuple_construct,
const opencl::program& p, const char* fn,
const opencl::spawn_config& cfg,
map_in_fun map_input, map_out_fun map_output,
Tuple&& xs,
detail::int_list<Is...>) const {
return actor_cast<actor>(impl::create(p, fn, cfg,std::move(map_input),
std::move(map_output),
std::move(std::get<Is>(xs))...));
}
};
template <class Signature>
struct cl_spawn_helper_deprecated;
template <class R, class... Ts>
struct cl_spawn_helper_deprecated<R(Ts...)> {
using input_types
= detail::type_list<typename opencl::carr_to_vec<Ts>::type...>;
using wrapped_input_types
= typename detail::tl_map<input_types,opencl::to_input_arg>::type;
using wrapped_output_type
= opencl::out<typename opencl::carr_to_vec<R>::type>;
using all_types =
typename detail::tl_push_back<
wrapped_input_types,
wrapped_output_type
>::type;
using values = typename detail::tl_apply<all_types, std::tuple>::type;
using helper_type =
typename detail::tl_apply<all_types, cl_spawn_helper>::type;
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& conf,
size_t result_size) const {
values xs;
std::get<tl_size<all_types>::value - 1>(xs) =
wrapped_output_type{[result_size](Ts&...){ return result_size; }};
helper_type f;
return f(tuple_construct{}, p, fn, conf, std::move(xs),
detail::get_indices(xs));
}
actor operator()(const opencl::program& p, const char* fn,
const opencl::spawn_config& conf,
size_t result_size,
typename helper_type::map_in_fun map_args,
typename helper_type::map_out_fun map_results) const {
values xs;
std::get<tl_size<all_types>::value - 1>(xs) =
wrapped_output_type{[result_size](Ts&...){ return result_size; }};
auto indices = detail::get_indices(xs);
helper_type f;
return f(tuple_construct{}, p, fn, conf,
std::move(map_args), std::move(map_results),
std::move(xs), std::move(indices));
}
};
} // namespace detail
/// 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_cl(const opencl::program& prog,
const char* fname,
const opencl::spawn_config& config,
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(prog, fname, config, std::move(x), std::move(xs)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template <class T, class... Ts>
typename std::enable_if<
opencl::is_opencl_arg<T>::value,
actor
>::type
spawn_cl(const char* source,
const char* fname,
const opencl::spawn_config& config,
T x,
Ts... xs) {
detail::cl_spawn_helper<T, Ts...> f;
return f(opencl::program::create(source), fname, config,
std::move(x), std::move(xs)...);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template <class Fun, class... Ts>
actor spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::spawn_config& config,
std::function<maybe<message> (message&)> map_args,
Fun map_result,
Ts... xs) {
detail::cl_spawn_helper<Ts...> f;
return f(prog, fname, config, std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template <class Fun, class... Ts>
actor spawn_cl(const char* source,
const char* fname,
const opencl::spawn_config& config,
std::function<maybe<message> (message&)> map_args,
Fun map_result,
Ts... xs) {
detail::cl_spawn_helper<Ts...> f;
return f(opencl::program::create(source), fname, config,
std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
// !!! Below are the deprecated spawn_cl functions !!!
// Signature first to mark them deprcated, gcc will complain otherwise
template <class Signature>
actor spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) CAF_DEPRECATED;
template <class Signature>
actor spawn_cl(const char* source,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) CAF_DEPRECATED;
template <class Signature, class Fun>
actor spawn_cl(const opencl::program& prog,
const char* fname,
std::function<maybe<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) CAF_DEPRECATED;
template <class Signature, class Fun>
actor spawn_cl(const char* source,
const char* fname,
std::function<maybe<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) CAF_DEPRECATED;
// now the implementations
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template <class Signature>
actor spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(prog, fname, opencl::spawn_config{dims, offset, local_dims},
result_size);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// 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 Signature>
actor spawn_cl(const char* source,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(opencl::program::create(source), fname,
opencl::spawn_config{dims, offset, local_dims}, result_size);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template <class Signature, class Fun>
actor spawn_cl(const opencl::program& prog,
const char* fname,
std::function<maybe<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(prog, fname, opencl::spawn_config{dims, offset, local_dims},
result_size, std::move(map_args), std::move(map_result));
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// 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 Signature, class Fun>
actor spawn_cl(const char* source,
const char* fname,
std::function<maybe<message> (message&)> map_args,
Fun map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset,
const opencl::dim_vec& local_dims,
size_t result_size) {
detail::cl_spawn_helper_deprecated<Signature> f;
return f(opencl::program::create(source), fname,
opencl::spawn_config{dims, offset, local_dims},
result_size, std::move(map_args), std::move(map_result));
}
} // namespace caf
#endif // CAF_SPAWN_CL_HPP
......@@ -24,13 +24,15 @@
#include <iostream>
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/all.hpp"
using namespace std;
using namespace caf;
using namespace caf::opencl;
using detail::limited_vector;
using caf::detail::limited_vector;
namespace {
......@@ -61,11 +63,16 @@ constexpr const char* kernel_source = R"__(
template<size_t Size>
class square_matrix {
public:
using value_type = fvec::value_type;
static constexpr size_t num_elements = Size * Size;
// allows serialization
template <class IO>
friend void serialize(IO& in_or_out, square_matrix& m, const unsigned int) {
in_or_out & m.data_;
}
square_matrix(square_matrix&&) = default;
square_matrix(const square_matrix&) = default;
square_matrix& operator=(square_matrix&&) = default;
......@@ -101,9 +108,7 @@ public:
const fvec& data() const { return data_; }
private:
fvec data_;
};
template<size_t Size>
......@@ -135,6 +140,7 @@ inline bool operator!=(const square_matrix<Size>& lhs,
using matrix_type = square_matrix<matrix_size>;
void multiplier(event_based_actor* self) {
auto& mngr = self->system().opencl_manager();
// create two matrices with ascending values
matrix_type m1;
......@@ -145,7 +151,7 @@ void multiplier(event_based_actor* self) {
cout << "calculating square of matrix:" << endl
<< to_string(m1) << endl;
auto unbox_args = [](message& msg) -> maybe<message> {
auto unbox_args = [](message& msg) -> optional<message> {
return msg.apply(
[](matrix_type& lhs, matrix_type& rhs) {
return make_message(std::move(lhs.data()), std::move(rhs.data()));
......@@ -171,14 +177,14 @@ void multiplier(event_based_actor* self) {
// used as response message
// from 6 : a description of the kernel signature using in/out/in_out classes
// with the argument type packed in vectors
auto worker = spawn_cl(kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
unbox_args, box_res,
in<fvec>{}, in<fvec>{}, out<fvec>{});
auto worker = mngr.spawn(kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
unbox_args, box_res,
in<fvec>{}, in<fvec>{}, out<fvec>{});
// send both matrices to the actor and
// wait for results in form of a matrix_type
self->sync_send(worker, move(m1), move(m2)).then(
self->request(worker, chrono::seconds(5), move(m1), move(m2)).then(
[](const matrix_type& result) {
cout << "result:" << endl << to_string(result);
}
......@@ -188,9 +194,12 @@ void multiplier(event_based_actor* self) {
int main() {
// matrix_type ist not a simple type,
// it must be annouced to libcaf
announce<matrix_type>("matrix_type");
spawn(multiplier);
await_all_actors_done();
shutdown();
actor_system_config cfg;
cfg.load<opencl::manager>()
.add_message_type<fvec>("float_vector")
.add_message_type<matrix_type>("square_matrix");
actor_system system{cfg};
system.spawn(multiplier);
system.await_all_actors_done();
return 0;
}
......@@ -23,13 +23,13 @@
#include <iostream>
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
#include "caf/opencl/all.hpp"
using namespace std;
using namespace caf;
using namespace caf::opencl;
using detail::limited_vector;
using caf::detail::limited_vector;
namespace {
......@@ -93,11 +93,13 @@ void multiplier(event_based_actor* self) {
// - local dimensions (optional)
// 4th to Nth arg: the kernel signature described by in/out/in_out classes
// that contain the argument type in their template, requires vectors
auto worker = spawn_cl(kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
in<fvec>{}, in<fvec>{}, out<fvec>{});
auto worker = self->system().opencl_manager().spawn(
kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
in<fvec>{}, in<fvec>{}, out<fvec>{}
);
// send both matrices to the actor and wait for a result
self->sync_send(worker, move(m1), move(m2)).then(
self->request(worker, chrono::seconds(5), move(m1), move(m2)).then(
[](const fvec& result) {
cout << "result: " << endl;
print_as_matrix(result);
......@@ -106,9 +108,11 @@ void multiplier(event_based_actor* self) {
}
int main() {
announce<fvec>("float_vector");
spawn(multiplier);
await_all_actors_done();
shutdown();
actor_system_config cfg;
cfg.load<opencl::manager>()
.add_message_type<fvec>("float_vector");
actor_system system{cfg};
system.spawn(multiplier);
system.await_all_actors_done();
return 0;
}
......@@ -19,6 +19,7 @@
#include <iostream>
#include "caf/logger.hpp"
#include "caf/string_algorithms.hpp"
#include "caf/opencl/global.hpp"
......@@ -31,8 +32,7 @@ namespace caf {
namespace opencl {
device device::create(context_ptr context, device_ptr device_id, unsigned id) {
CAF_LOGF_DEBUG("creating device for opencl device id '"
<< device_id.get() << "' with id '" << id << "'");
CAF_LOG_DEBUG("creating device for opencl device with id:" << CAF_ARG(id));
// look up properties we need to create the command queue
auto supported = info<cl_ulong>(device_id, CL_DEVICE_QUEUE_PROPERTIES);
bool profiling = supported & CL_QUEUE_PROFILING_ENABLE;
......
......@@ -18,26 +18,19 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include "caf/detail/type_list.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/metainfo.hpp"
#include "caf/opencl/manager.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
using namespace std;
namespace caf {
namespace opencl {
metainfo* metainfo::instance() {
auto sid = detail::singletons::opencl_plugin_id;
auto fac = [] { return new metainfo; };
auto res = detail::singletons::get_plugin_singleton(sid, fac);
return static_cast<metainfo*>(res);
}
const std::vector<device>& metainfo::get_devices() const {
return platforms_.front().get_devices();
}
const maybe<const device&> metainfo::get_device(size_t id) const{
const optional<const device&> manager::get_device(size_t id) const {
if (platforms_.empty())
return none;
size_t to = 0;
......@@ -50,7 +43,7 @@ const maybe<const device&> metainfo::get_device(size_t id) const{
return none;
}
void metainfo::initialize() {
void manager::init(actor_system_config&) {
// get number of available platforms
auto num_platforms = v1get<cl_uint>(CAF_CLF(clGetPlatformIDs));
// get platform ids
......@@ -67,11 +60,114 @@ void metainfo::initialize() {
}
}
void metainfo::dispose() {
delete this;
void manager::start() {
// nop
}
void manager::stop() {
// nop
}
actor_system::module::id_t manager::id() const {
return actor_system::module::opencl_manager;
}
void* manager::subtype_ptr() {
return this;
}
actor_system::module* manager::make(actor_system& sys,
caf::detail::type_list<>) {
return new manager{sys};
}
program manager::create_program(const char* kernel_source, const char* options,
uint32_t device_id) {
auto dev = get_device(device_id);
if (! dev) {
ostringstream oss;
oss << "No device with id '" << device_id << "' found.";
CAF_LOG_ERROR(CAF_ARG(oss.str()));
throw runtime_error(oss.str());
}
return create_program(kernel_source, options, *dev);
}
program manager::create_program(const char* kernel_source, const char* options,
const device& dev) {
// create program object from kernel source
size_t kernel_source_length = strlen(kernel_source);
program_ptr pptr;
auto rawptr = v2get(CAF_CLF(clCreateProgramWithSource), dev.context_.get(),
cl_uint{1}, &kernel_source, &kernel_source_length);
pptr.reset(rawptr, false);
// build programm from program object
auto dev_tmp = dev.device_id_.get();
cl_int err = clBuildProgram(pptr.get(), 1, &dev_tmp,
options, nullptr, nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
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) {
size_t buildlog_buffer_size = 0;
// get the log length
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buildlog_buffer_size), nullptr,
&buildlog_buffer_size);
vector<char> buffer(buildlog_buffer_size);
// fill the buffer with buildlog informations
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buffer[0]) * buildlog_buffer_size,
buffer.data(), nullptr);
ostringstream ss;
ss << "Build log:\n" << string(buffer.data())
<< "\n#######################################";
CAF_LOG_ERROR(CAF_ARG(ss.str()));
}
#endif
throw runtime_error(oss.str());
}
map<string, kernel_ptr> available_kernels;
cl_uint number_of_kernels = 0;
clCreateKernelsInProgram(pptr.get(), 0, nullptr, &number_of_kernels);
vector<cl_kernel> kernels(number_of_kernels);
err = clCreateKernelsInProgram(pptr.get(), number_of_kernels, kernels.data(),
nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clCreateKernelsInProgram: " << get_opencl_error(err);
throw runtime_error(oss.str());
} else {
for (cl_uint i = 0; i < number_of_kernels; ++i) {
size_t ret_size;
clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, 0, nullptr, &ret_size);
vector<char> name(ret_size);
err = clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, ret_size,
reinterpret_cast<void*>(name.data()), nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetKernelInfo (CL_KERNEL_FUNCTION_NAME): "
<< get_opencl_error(err);
throw runtime_error(oss.str());
}
kernel_ptr kernel;
kernel.reset(move(kernels[i]));
available_kernels.emplace(string(name.data()), move(kernel));
}
}
return {dev.context_, dev.command_queue_, pptr, move(available_kernels)};
}
manager::manager(actor_system& sys) : system_(sys){
// nop
}
void metainfo::stop() {
manager::~manager() {
// nop
}
......
......@@ -33,9 +33,10 @@ void throwcl(const char* fname, cl_int err) {
}
void pfn_notify(const char* errinfo, const void*, size_t, void*) {
CAF_LOGF_ERROR("\n##### Error message via pfn_notify #####\n"
<< errinfo <<
"\n########################################");
CAF_LOG_ERROR("\n##### Error message via pfn_notify #####\n"
<< errinfo <<
"\n########################################");
static_cast<void>(errinfo); // remove warning
}
} // namespace opencl
......
......@@ -63,7 +63,7 @@ platform platform::create(cl_platform_id platform_id,
}
if (device_information.empty()) {
string errstr = "no devices for the platform found";
CAF_LOGF_ERROR(errstr);
CAF_LOG_ERROR(CAF_ARG(errstr));
throw runtime_error(move(errstr));
}
auto name = platform_info(platform_id, CL_PLATFORM_NAME);
......
......@@ -23,8 +23,8 @@
#include <cstring>
#include <iostream>
#include "caf/opencl/manager.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/metainfo.hpp"
#include "caf/opencl/opencl_err.hpp"
using namespace std;
......@@ -32,63 +32,13 @@ using namespace std;
namespace caf {
namespace opencl {
program::program(context_ptr context, command_queue_ptr queue, program_ptr prog)
program::program(context_ptr context, command_queue_ptr queue, program_ptr prog,
map<string, kernel_ptr> available_kernels)
: context_(move(context)),
program_(move(prog)),
queue_(move(queue)) {}
program program::create(const char* kernel_source, const char* options,
uint32_t device_id) {
auto info = metainfo::instance();
auto dev = info->get_device(device_id);
if (! dev) {
ostringstream oss;
oss << "No device with id '" << device_id << "' found.";
CAF_LOGF_ERROR(oss.str());
throw runtime_error(oss.str());
}
return program::create(kernel_source, options, *dev);
}
program program::create(const char* kernel_source, const char* options,
const device& dev) {
// create program object from kernel source
size_t kernel_source_length = strlen(kernel_source);
program_ptr pptr;
auto rawptr = v2get(CAF_CLF(clCreateProgramWithSource), dev.context_.get(),
cl_uint{1}, &kernel_source, &kernel_source_length);
pptr.reset(rawptr, false);
// build programm from program object
auto dev_tmp = dev.device_id_.get();
cl_int err = clBuildProgram(pptr.get(), 1, &dev_tmp,
options, nullptr, nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clBuildProgram: " << get_opencl_error(err);
// the build log will be printed by the
// pfn_notify (see metainfo.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) {
size_t buildlog_buffer_size = 0;
// get the log length
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buildlog_buffer_size), nullptr,
&buildlog_buffer_size);
vector<char> buffer(buildlog_buffer_size);
// fill the buffer with buildlog informations
clGetProgramBuildInfo(pptr.get(), dev_tmp, CL_PROGRAM_BUILD_LOG,
sizeof(buffer[0]) * buildlog_buffer_size,
buffer.data(), nullptr);
CAF_LOGF_ERROR("Build log:\n" + string(buffer.data()) +
"\n########################################");
}
#endif
throw runtime_error(oss.str());
}
return {dev.context_, dev.command_queue_, pptr};
queue_(move(queue)),
available_kernels_(move(available_kernels)) {
// nop
}
} // namespace opencl
......
......@@ -8,13 +8,15 @@
#include <algorithm>
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
#include "caf/system_messages.hpp"
#include "caf/opencl/all.hpp"
using namespace std;
using namespace caf;
using namespace caf::opencl;
using detail::limited_vector;
using caf::detail::limited_vector;
namespace {
......@@ -108,10 +110,13 @@ constexpr const char* kernel_source_inout = R"__(
template<size_t Size>
class square_matrix {
public:
using value_type = ivec::value_type;
static constexpr size_t num_elements = Size * Size;
static void announce() {
caf::announce<square_matrix>("square_matrix", &square_matrix::data_);
// allows serialization
template <class IO>
friend void serialize(IO& in_or_out, square_matrix& m, const unsigned int) {
in_or_out & m.data_;
}
square_matrix(square_matrix&&) = default;
......@@ -198,7 +203,7 @@ void check_vector_results(const string& description,
auto cond = (expected == result);
CAF_CHECK(cond);
if (! cond) {
CAF_TEST_INFO(description << " failed.");
CAF_ERROR(description << " failed.");
cout << "Expected: " << endl;
for (size_t i = 0; i < expected.size(); ++i) {
cout << " " << expected[i];
......@@ -211,48 +216,47 @@ void check_vector_results(const string& description,
}
}
void test_opencl() {
auto info = metainfo::instance();
auto opt = info->get_device_if([](const device&){ return true; });
void test_opencl(actor_system& sys) {
static_cast<void>(sys);
auto& mngr = sys.opencl_manager();
auto opt = mngr.get_device_if([](const device&){ return true; });
if (! opt)
CAF_TEST_ERROR("No OpenCL device found.");
CAF_ERROR("No OpenCL device found.");
auto dev = *opt;
scoped_actor self;
scoped_actor self{sys};
self->set_default_handler(
[=](local_actor*, const type_erased_tuple* x) -> result<message> {
CAF_ERROR("unexpected message" << x->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 = spawn_cl(program::create(kernel_source, "", dev), kernel_name,
opencl::spawn_config{dims{matrix_size, matrix_size}},
opencl::in<ivec>{}, opencl::out<ivec>{});
auto w1 = mngr.spawn(mngr.create_program(kernel_source, "", dev), kernel_name,
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("Simple matrix multiplication using vectors"
"(kernel wrapped in program)",
expected1, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
opencl::spawn_config cfg2{dims{matrix_size, matrix_size}};
auto w2 = spawn_cl(kernel_source, kernel_name, cfg2,
opencl::in<ivec>{}, opencl::out<ivec>{});
auto w2 = mngr.spawn(kernel_source, kernel_name, cfg2,
opencl::in<int*>{}, opencl::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",
expected1, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
const matrix_type expected2(move(expected1));
auto map_arg = [](message& msg) -> maybe<message> {
auto map_arg = [](message& msg) -> optional<message> {
return msg.apply(
[](matrix_type& mx) {
return make_message(move(mx.data()));
......@@ -263,67 +267,55 @@ void test_opencl() {
return make_message(matrix_type{move(result)});
};
opencl::spawn_config cfg3{dims{matrix_size, matrix_size}};
auto w3 = spawn_cl(program::create(kernel_source), kernel_name, cfg3,
map_arg, map_res,
opencl::in<ivec>{}, opencl::out<ivec>{});
auto w3 = mngr.spawn(mngr.create_program(kernel_source), kernel_name, cfg3,
map_arg, map_res,
opencl::in<int*>{}, opencl::out<int*>{});
self->send(w3, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type "
"(kernel wrapped in program)",
expected2.data(), result.data());
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
opencl::spawn_config cfg4{dims{matrix_size, matrix_size}};
auto w4 = spawn_cl(kernel_source, kernel_name, cfg4,
map_arg, map_res,
opencl::in<ivec>{}, opencl::out<ivec>{});
auto w4 = mngr.spawn(kernel_source, kernel_name, cfg4,
map_arg, map_res,
opencl::in<int*>{}, opencl::out<int*>{});
self->send(w4, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type",
expected2.data(), result.data());
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
CAF_TEST_INFO("Expecting exception (compiling invalid kernel, "
"semicolon is missing).");
CAF_MESSAGE("Expecting exception (compiling invalid kernel, "
"semicolon is missing).");
try {
auto create_error = program::create(kernel_source_error);
auto create_error = mngr.create_program(kernel_source_error);
}
catch (const exception& exc) {
auto cond = (strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE",
exc.what()) == 0);
CAF_CHECK(cond);
if (! cond) {
CAF_TEST_INFO("Wrong exception cought for program build failure.");
CAF_ERROR("Wrong exception cought for program build failure.");
}
}
// test for opencl compiler flags
auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag);
auto prog5 = mngr.create_program(kernel_source_compiler_flag, compiler_flag);
opencl::spawn_config cfg5{dims{array_size}};
auto w5 = spawn_cl(prog5, kernel_name_compiler_flag, cfg5,
opencl::in<ivec>{}, opencl::out<ivec>{});
auto w5 = mngr.spawn(prog5, kernel_name_compiler_flag, cfg5,
opencl::in<int*>{}, opencl::out<int*>{});
self->send(w5, make_iota_vector<int>(array_size));
auto expected3 = make_iota_vector<int>(array_size);
self->receive (
[&](const ivec& result) {
check_vector_results("Passing compiler flags", expected3, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
auto dev6 = info->get_device_if([](const device& d) {
auto dev6 = mngr.get_device_if([](const device& d) {
return d.get_device_type() != cpu;
});
if (dev6) {
......@@ -338,14 +330,14 @@ void test_opencl() {
int n = static_cast<int>(arr6.capacity());
generate(arr6.begin(), arr6.end(), [&]{ return --n; });
opencl::spawn_config cfg6{dims{reduce_global_size},
dims{ /* no offsets */ },
dims{ }, // no offset
dims{reduce_local_size}};
auto get_result_size_6 = [reduce_result_size](const ivec&) {
return reduce_result_size;
};
auto w6 = spawn_cl(program::create(kernel_source_reduce, "", *dev6),
kernel_name_reduce, cfg6,
opencl::in<ivec>{}, opencl::out<ivec>{get_result_size_6});
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,
......@@ -354,10 +346,6 @@ void test_opencl() {
self->receive(
[&](const ivec& result) {
check_vector_results("Passing size for the output", expected4, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
}
......@@ -367,20 +355,16 @@ void test_opencl() {
};
// constant memory arguments
const ivec arr7{problem_size};
auto w7 = spawn_cl(kernel_source_const, kernel_name_const,
opencl::spawn_config{dims{problem_size}},
opencl::in<ivec>{},
opencl::out<ivec>{get_result_size_7});
auto w7 = mngr.spawn(kernel_source_const, kernel_name_const,
opencl::spawn_config{dims{problem_size}},
opencl::in<ivec>{},
opencl::out<ivec>{get_result_size_7});
self->send(w7, move(arr7));
ivec expected5(problem_size);
fill(begin(expected5), end(expected5), static_cast<int>(problem_size));
self->receive(
[&](const ivec& result) {
check_vector_results("Using const input argument", expected5, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
// test in_out argument type
......@@ -388,25 +372,23 @@ void test_opencl() {
ivec expected9{input9};
transform(begin(expected9), end(expected9), begin(expected9),
[](const int& val){ return val * 2; });
auto w9 = spawn_cl(kernel_source_inout, kernel_name_inout,
spawn_config{dims{problem_size}},
opencl::in_out<ivec>{});
auto w9 = mngr.spawn(kernel_source_inout, kernel_name_inout,
spawn_config{dims{problem_size}},
opencl::in_out<ivec>{});
self->send(w9, move(input9));
self->receive(
[&](const ivec& result) {
check_vector_results("Testing in_out arugment", expected9, result);
},
others >> [&] {
CAF_TEST_ERROR("Unexpected message "
<< to_string(self->current_message()));
}
);
}
CAF_TEST(test_opencl) {
announce<ivec>("ivec");
matrix_type::announce();
test_opencl();
await_all_actors_done();
shutdown();
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_opencl(system);
system.await_all_actors_done();
}
#define CAF_SUITE opencl
#include "caf/test/unit_test.hpp"
#include <vector>
#include <iomanip>
#include <cassert>
#include <iostream>
#include <algorithm>
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
CAF_PUSH_NO_DEPRECATED_WARNING
using namespace caf;
using namespace caf::opencl;
using detail::limited_vector;
namespace {
using ivec = std::vector<int>;
constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32;
constexpr int magic_number = 23;
constexpr const char* kernel_name = "matrix_square";
constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_const = "const_mod";
constexpr const char* compiler_flag = "-D CAF_OPENCL_TEST_FLAG";
constexpr const char* kernel_source = R"__(
__kernel void matrix_square(__global int* matrix,
__global int* output) {
size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0);
size_t y = get_global_id(1);
int result = 0;
for (size_t idx = 0; idx < size; ++idx) {
result += matrix[idx + y * size] * matrix[x + idx * size];
}
output[x + y * size] = result;
}
)__";
constexpr const char* kernel_source_error = R"__(
__kernel void missing(__global int*) {
size_t semicolon
}
)__";
constexpr const char* kernel_source_compiler_flag = R"__(
__kernel void compiler_flag(__global int* input,
__global int* output) {
size_t x = get_global_id(0);
# ifdef CAF_OPENCL_TEST_FLAG
output[x] = input[x];
# else
output[x] = 0;
# endif
}
)__";
constexpr const char* kernel_source_const = R"__(
__kernel void const_mod(__constant int* input,
__global int* output) {
size_t idx = get_global_id(0);
output[idx] = input[0];
}
)__";
} // namespace <anonymous>
template<size_t Size>
class square_matrix {
public:
static constexpr size_t num_elements = Size * Size;
static void announce() {
caf::announce<square_matrix>("square_matrix", &square_matrix::data_);
}
square_matrix(square_matrix&&) = default;
square_matrix(const square_matrix&) = default;
square_matrix& operator=(square_matrix&&) = default;
square_matrix& operator=(const square_matrix&) = default;
square_matrix() : data_(num_elements) {
// nop
}
explicit square_matrix(ivec d) : data_(move(d)) {
assert(data_.size() == num_elements);
}
float& operator()(size_t column, size_t row) {
return data_[column + row * Size];
}
const float& operator()(size_t column, size_t row) const {
return data_[column + row * Size];
}
typedef typename ivec::const_iterator const_iterator;
const_iterator begin() const {
return data_.begin();
}
const_iterator end() const {
return data_.end();
}
ivec& data() {
return data_;
}
const ivec& data() const {
return data_;
}
void data(ivec new_data) {
data_ = std::move(new_data);
}
private:
ivec data_;
};
template <class T>
std::vector<T> make_iota_vector(size_t num_elements) {
std::vector<T> result;
result.resize(num_elements);
std::iota(result.begin(), result.end(), T{0});
return result;
}
template <size_t Size>
square_matrix<Size> make_iota_matrix() {
square_matrix<Size> result;
std::iota(result.data().begin(), result.data().end(), 0);
return result;
}
template<size_t Size>
bool operator==(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) {
return lhs.data() == rhs.data();
}
template<size_t Size>
bool operator!=(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) {
return ! (lhs == rhs);
}
using matrix_type = square_matrix<matrix_size>;
template <class T>
void check_vector_results(const std::string& description,
const std::vector<T>& expected,
const std::vector<T>& result) {
auto cond = (expected == result);
CAF_CHECK(cond);
if (! cond) {
CAF_TEST_INFO(description << " failed.");
std::cout << "Expected: " << std::endl;
for (size_t i = 0; i < expected.size(); ++i) {
std::cout << " " << expected[i];
}
std::cout << std::endl << "Received: " << std::endl;
for (size_t i = 0; i < result.size(); ++i) {
std::cout << " " << result[i];
}
std::cout << std::endl;
}
}
void test_opencl_deprecated() {
scoped_actor self;
const ivec expected1{ 56, 62, 68, 74,
152, 174, 196, 218,
248, 286, 324, 362,
344, 398, 452, 506};
auto w1 = spawn_cl<ivec (ivec)>(program::create(kernel_source),
kernel_name,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
self->receive (
[&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors"
"(kernel wrapped in program)",result, expected1);
}
);
auto w2 = spawn_cl<ivec (ivec)>(kernel_source, kernel_name,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w2, make_iota_vector<int>(matrix_size * matrix_size));
self->receive (
[&](const ivec& result) {
check_vector_results("Simple matrix multiplication using vectors",
expected1, result);
}
);
const matrix_type expected2(std::move(expected1));
auto map_arg = [](message& msg) -> maybe<message> {
return msg.apply(
[](matrix_type& mx) {
return make_message(std::move(mx.data()));
}
);
};
auto map_res = [](ivec& result) -> message {
return make_message(matrix_type{std::move(result)});
};
auto w3 = spawn_cl<ivec (ivec)>(program::create(kernel_source),
kernel_name, map_arg, map_res,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w3, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type "
"(kernel wrapped in program)",
expected2.data(), result.data());
}
);
auto w4 = spawn_cl<ivec (ivec)>(kernel_source, kernel_name,
map_arg, map_res,
limited_vector<size_t, 3>{matrix_size,
matrix_size});
self->send(w4, make_iota_matrix<matrix_size>());
self->receive (
[&](const matrix_type& result) {
check_vector_results("Matrix multiplication with user defined type",
expected2.data(), result.data());
}
);
CAF_TEST_INFO("Expecting exception (compiling invalid kernel, "
"semicolon is missing).");
try {
auto create_error = program::create(kernel_source_error);
}
catch (const std::exception& exc) {
auto cond = (strcmp("clBuildProgram: CL_BUILD_PROGRAM_FAILURE",
exc.what()) == 0);
CAF_CHECK(cond);
if (! cond) {
CAF_TEST_INFO("Wrong exception cought for program build failure.");
}
}
// test for opencl compiler flags
auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag);
auto w5 = spawn_cl<ivec (ivec)>(prog5, kernel_name_compiler_flag,
limited_vector<size_t, 3>{array_size});
self->send(w5, make_iota_vector<int>(array_size));
auto expected3 = make_iota_vector<int>(array_size);
self->receive (
[&](const ivec& result) {
check_vector_results("Passing compiler flags", expected3, result);
}
);
// constant memory arguments
const ivec arr7{magic_number};
auto w7 = spawn_cl<ivec (ivec)>(kernel_source_const, kernel_name_const,
limited_vector<size_t, 3>{magic_number});
self->send(w7, move(arr7));
ivec expected5(magic_number);
fill(begin(expected5), end(expected5), magic_number);
self->receive(
[&](const ivec& result) {
check_vector_results("Using const input argument", expected5, result);
}
);
}
CAF_TEST(test_opencl_deprecated) {
std::cout << "Staring deprecated OpenCL test" << std::endl;
announce<ivec>("ivec");
matrix_type::announce();
test_opencl_deprecated();
await_all_actors_done();
shutdown();
std::cout << "Done with depreacted OpenCL test" << std::endl;
}
CAF_POP_WARNINGS
Markdown is supported
0%
or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment