Commit 9b3c07ad authored by Dominik Charousset's avatar Dominik Charousset

Merge pull request #581

parents e520075d 2b699267
[submodule "libcaf_opencl"]
path = libcaf_opencl
url = ../opencl.git
[submodule "libcaf_python/third_party/pybind"]
path = libcaf_python/third_party/pybind
url = https://github.com/pybind/pybind11.git
Subproject commit cc4e4ef5ad92e3ee78e2f2897027e5f48cbc8e93
cmake_minimum_required(VERSION 2.8)
project(caf_opencl C CXX)
# get header files; only needed by CMake generators,
# e.g., for creating proper Xcode projects
file(GLOB_RECURSE LIBCAF_OPENCL_HDRS "caf/*.hpp")
add_custom_target(libcaf_opencl)
# list cpp files excluding platform-dependent files
set (LIBCAF_OPENCL_SRCS
src/global.cpp
src/manager.cpp
src/program.cpp
src/opencl_err.cpp
src/platform.cpp
src/device.cpp)
# build shared library if not compiling static only
if(NOT CAF_BUILD_STATIC_ONLY)
add_library(libcaf_opencl_shared SHARED ${LIBCAF_OPENCL_SRCS}
${LIBCAF_OPENCL_HDRS} ${OpenCL_INCLUDE_DIRS})
target_link_libraries(libcaf_opencl_shared ${LD_FLAGS}
${CAF_LIBRARY_CORE}
${OpenCL_LIBRARIES})
set_target_properties(libcaf_opencl_shared
PROPERTIES
SOVERSION "${CAF_VERSION}"
VERSION "${CAF_VERSION}"
OUTPUT_NAME caf_opencl)
if(NOT WIN32)
install(TARGETS libcaf_opencl_shared LIBRARY DESTINATION lib)
endif()
endif()
# build static library only if --build-static or --build-static-only was set
if(CAF_BUILD_STATIC_ONLY OR CAF_BUILD_STATIC)
add_library(libcaf_opencl_static STATIC ${LIBCAF_OPENCL_HDRS} ${LIBCAF_OPENCL_SRCS})
target_link_libraries(libcaf_opencl_static ${LD_FLAGS}
${CAF_LIBRARY_CORE_STATIC}
${OpenCL_LIBRARIES})
set_target_properties(libcaf_opencl_static PROPERTIES OUTPUT_NAME caf_opencl_static)
install(TARGETS libcaf_opencl_static ARCHIVE DESTINATION lib)
endif()
link_directories(${LD_DIRS})
include_directories(. ${INCLUDE_DIRS})
# install includes
install(DIRECTORY caf/ DESTINATION include/caf FILES_MATCHING PATTERN "*.hpp")
OpenCL Actors
=============
This module eases the use of OpenCL with CAF. See our [Wiki page](https://github.com/actor-framework/actor-framework/wiki/OpenCL-Actors) for details.
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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/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/nd_range.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
#include "caf/opencl/detail/command_helper.hpp"
namespace caf {
namespace opencl {
class manager;
template <bool PassConfig, 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_input_type>::type;
using input_mapping = typename std::conditional<PassConfig,
std::function<optional<message> (nd_range&, 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 detail::output_function_sig<output_types>::type;
using processing_list = typename cl_arg_info_list<arg_types>::type;
using command_type = typename detail::command_sig<actor_facade, output_types>::type;
typename detail::il_indices<arg_types>::type indices;
using evnt_vec = std::vector<cl_event>;
using mem_vec = std::vector<detail::raw_mem_ptr>;
using len_vec = std::vector<size_t>;
using out_tup = typename detail::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 nd_range& range,
input_mapping map_args, output_mapping map_result,
Ts&&... xs) {
if (range.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() != range.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(range.offsets(), "offsets");
check_vec(range.local_dimensions(), "local dimensions");
auto& sys = actor_conf.host->system();
auto itr = prog->available_kernels_.find(kernel_name);
if (itr == prog->available_kernels_.end()) {
detail::raw_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_conf),
prog, kernel, range,
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_conf),
prog, itr->second, range,
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),
range_
);
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(std::move(sender), mid, {},
std::move(content)), host);
}
actor_facade(actor_config actor_conf, const program_ptr prog,
detail::raw_kernel_ptr kernel, nd_range range,
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_),
range_(std::move(range)),
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(range_.dimensions()),
std::end(range_.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 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_, detail::raw_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>
detail::enable_if_t<!Q, bool> 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>
detail::enable_if_t<Q, bool> map_arguments(message& content) {
if (map_args_) {
auto mapped = map_args_(range_, content);
if (!mapped) {
CAF_LOG_ERROR("Mapping argumentes failed.");
return false;
}
content = std::move(*mapped);
}
return true;
}
detail::raw_kernel_ptr kernel_;
detail::raw_program_ptr program_;
detail::raw_context_ptr context_;
detail::raw_command_queue_ptr queue_;
nd_range range_;
input_mapping map_args_;
output_mapping map_results_;
std::tuple<Ts...> kernel_signature_;
size_t default_length_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_ACTOR_FACADE_HPP
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* Raphael Hiesgen <raphael.hiesgen (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_ALL_HPP
#define CAF_OPENCL_ALL_HPP
#include "caf/opencl/manager.hpp"
#endif // CAF_OPENCL_ALL_HPP
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* Raphael Hiesgen <raphael.hiesgen (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_ARGUMENTS
#define CAF_OPENCL_ARGUMENTS
#include <functional>
#include <type_traits>
#include "caf/message.hpp"
#include "caf/optional.hpp"
#include "caf/opencl/mem_ref.hpp"
#include "caf/opencl/detail/core.hpp"
namespace caf {
namespace opencl {
namespace detail {
template <class T, class F>
std::function<optional<T> (message&)> res_or_none(F fun) {
return [fun](message& msg) -> optional<T> {
auto res = msg.apply(fun);
T result;
if (res) {
res->apply([&](size_t x) { result = x; });
return result;
}
return none;
};
}
template <class F, class T>
T try_apply_fun(F& fun, message& msg, const T& fallback) {
if (fun) {
auto res = fun(msg);
if (res)
return *res;
}
return fallback;
}
} // namespace detail
// 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
/// of work items at runtime.
struct dummy_size_calculator {
template <class... Ts>
size_t operator()(Ts&&...) const {
return 0;
}
};
/// Common parent for opencl argument tags for the spawn function.
struct arg_tag {};
/// Empty tag as an alternative for conditional inheritance.
struct empty_tag {};
/// Tags the argument as input which requires initialization through a message.
struct input_tag {};
/// Tags the argument as output which includes its buffer in the result message.
struct output_tag {};
/// Tags the argument to require specification of the size of its buffer.
struct requires_size_tag {};
/// Tags the argument as a reference and not a value.
struct is_ref_tag;
/// Mark a spawn argument as input only
template <class Arg, class Tag = val>
struct in : arg_tag, input_tag {
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 = detail::decay_t<Arg>;
};
/// Mark a spawn argument as input and output
template <class Arg, class TagIn = val, class TagOut = val>
struct in_out : arg_tag, input_tag, output_tag {
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 = detail::decay_t<Arg>;
};
/// Mark a spawn argument as output only
template <class Arg, class Tag = val>
struct out : arg_tag, output_tag, requires_size_tag {
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 = detail::decay_t<Arg>;
out() = default;
template <class F>
out(F fun) : fun_{detail::res_or_none<size_t>(fun)} {
// nop
}
optional<size_t> operator()(message& msg) const {
return detail::try_apply_fun(fun_, msg, 0UL);
}
std::function<optional<size_t> (message&)> fun_;
};
/// Mark a spawn argument as on-device scratch space
template <class Arg>
struct scratch : arg_tag, requires_size_tag {
using arg_type = detail::decay_t<Arg>;
scratch() = default;
template <class F>
scratch(F fun) : fun_{detail::res_or_none<size_t>(fun)} {
// nop
}
optional<size_t> operator()(message& msg) const {
return detail::try_apply_fun(fun_, msg, 0UL);
}
std::function<optional<size_t> (message&)> fun_;
};
/// Mark a spawn argument as a local memory argument. This argument cannot be
/// initalized from the CPU, but requires specification of its size. An
/// optional function allows calculation of the size depeding on the input
/// message.
template <class Arg>
struct local : arg_tag, requires_size_tag {
using arg_type = detail::decay_t<Arg>;
local() = default;
local(size_t size) : size_(size) { }
template <class F>
local(size_t size, F fun)
: size_(size), fun_{detail::res_or_none<size_t>(fun)} {
// nop
}
size_t operator()(message& msg) const {
return detail::try_apply_fun(fun_, msg, size_);
}
size_t size_;
std::function<optional<size_t> (message&)> fun_;
};
/// Mark a spawn argument as a private argument. Requires a default value but
/// can optionally be calculated depending on the input through a passed
/// function.
template <class Arg, class Tag = hidden>
struct priv : arg_tag, std::conditional<std::is_same<Tag, val>::value,
input_tag, empty_tag>::type {
static_assert(std::is_same<Tag, val>::value ||
std::is_same<Tag, hidden>::value,
"Argument of type `priv` must be either a value or hidden.");
using tag_type = Tag;
using arg_type = detail::decay_t<Arg>;
priv() = default;
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 tagged as hidden.");
}
template <class F>
priv(Arg val, F fun) : value_(val), fun_{detail::res_or_none<Arg>(fun)} {
static_assert(std::is_same<Tag, hidden>::value,
"Argument of type `priv` can only be initialized with a value"
" if it is tagged as hidden.");
}
Arg operator()(message& msg) const {
return detail::try_apply_fun(fun_, msg, value_);
}
Arg value_;
std::function<optional<Arg> (message&)> fun_;
};
///Cconverts C arrays, i.e., pointers, to vectors.
template <class T>
struct carr_to_vec {
using type = T;
};
template <class T>
struct carr_to_vec<T*> {
using type = std::vector<T>;
};
/// Filter types for any argument type.
template <class T>
struct is_opencl_arg : std::is_base_of<arg_tag, T> {};
/// Filter type lists for input arguments
template <class T>
struct is_input_arg : std::is_base_of<input_tag, T> {};
/// Filter type lists for output arguments
template <class T>
struct is_output_arg : std::is_base_of<output_tag, T> {};
/// Filter for arguments that require size
template <class T>
struct requires_size_arg : std::is_base_of<requires_size_tag, T> {};
/// Filter mem_refs
template <class T>
struct is_ref_type : std::is_base_of<is_ref_tag, T> {};
template <class T>
struct is_val_type
: std::integral_constant<bool, !std::is_base_of<is_ref_tag, T>::value> {};
/// extract types
template <class T>
struct extract_type { };
template <class T, class Tag>
struct extract_type<in<T, Tag>> {
using type = detail::decay_t<typename carr_to_vec<T>::type>;
};
template <class T, class TagIn, class TagOut>
struct extract_type<in_out<T, TagIn, TagOut>> {
using type = detail::decay_t<typename carr_to_vec<T>::type>;
};
template <class T, class Tag>
struct extract_type<out<T, Tag>> {
using type = detail::decay_t<typename carr_to_vec<T>::type>;
};
template <class T>
struct extract_type<scratch<T>> {
using type = detail::decay_t<typename carr_to_vec<T>::type>;
};
template <class T>
struct extract_type<local<T>> {
using type = detail::decay_t<typename carr_to_vec<T>::type>;
};
template <class T, class Tag>
struct extract_type<priv<T, Tag>> {
using type = detail::decay_t<typename carr_to_vec<T>::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
struct message_from_results {
template <class T, class... Ts>
message operator()(T& x, Ts&... xs) {
return make_message(std::move(x), std::move(xs)...);
}
template <class... Ts>
message operator()(std::tuple<Ts...>& values) {
return apply_args(*this, detail::get_indices(values), values);
}
};
/// 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
template <class T>
struct to_input_arg {
using type = in<T>;
};
template <class T>
struct to_output_arg {
using type = out<T>;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_ARGUMENTS
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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_COMMAND_HPP
#define CAF_OPENCL_COMMAND_HPP
#include <tuple>
#include <vector>
#include <numeric>
#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/scope_guard.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/nd_range.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace caf {
namespace opencl {
/// 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 {
public:
using result_types = detail::type_list<Ts...>;
command(response_promise promise,
strong_actor_ptr parent,
std::vector<cl_event> events,
std::vector<detail::raw_mem_ptr> inputs,
std::vector<detail::raw_mem_ptr> outputs,
std::vector<detail::raw_mem_ptr> scratches,
std::vector<size_t> lengths,
message msg,
std::tuple<Ts...> output_tuple,
nd_range range)
: lengths_(std::move(lengths)),
promise_(std::move(promise)),
cl_actor_(std::move(parent)),
mem_in_events_(std::move(events)),
input_buffers_(std::move(inputs)),
output_buffers_(std::move(outputs)),
scratch_buffers_(std::move(scratches)),
results_(std::move(output_tuple)),
msg_(std::move(msg)),
range_(std::move(range)) {
// nop
}
~command() override {
for (auto& e : mem_in_events_) {
if (e)
v1callcl(CAF_CLF(clReleaseEvent), e);
}
for (auto& e : mem_out_events_) {
if (e)
v1callcl(CAF_CLF(clReleaseEvent), e);
}
}
/// 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>
detail::enable_if_t<!detail::tl_forall<Q, is_ref_type>::value>
enqueue() {
// Errors in this function can not be handled by opencl_err.hpp
// because they require non-standard error handling
CAF_LOG_TRACE("");
this->ref(); // reference held by the OpenCL comand 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_));
// OpenCL expects cl_uint (unsigned int), hence the cast
mem_out_events_.emplace_back();
auto success = invoke_cl(
clEnqueueNDRangeKernel, parent->queue_.get(), parent->kernel_.get(),
static_cast<unsigned int>(range_.dimensions().size()),
data_or_nullptr(range_.offsets()),
data_or_nullptr(range_.dimensions()),
data_or_nullptr(range_.local_dimensions()),
static_cast<unsigned int>(mem_in_events_.size()),
(mem_in_events_.empty() ? nullptr : mem_in_events_.data()),
&mem_out_events_.back()
);
if (!success)
return;
size_t pos = 0;
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__)
success = invoke_cl(clEnqueueMarkerWithWaitList, parent->queue_.get(),
static_cast<unsigned int>(mem_out_events_.size()),
mem_out_events_.data(), &marker_event);
#else
success = invoke_cl(clEnqueueMarker, parent->queue_.get(), &marker_event);
#endif
callback_.reset(marker_event, false);
if (!success)
return;
auto cb = [](cl_event, cl_int, void* data) {
auto cmd = reinterpret_cast<command*>(data);
cmd->handle_results();
cmd->deref();
};
if (!invoke_cl(clSetEventCallback, callback_.get(), CL_COMPLETE,
std::move(cb), this))
return;
if (clFlush(parent->queue_.get()) != CL_SUCCESS)
CAF_LOG_ERROR("clFlush: " << CAF_ARG(get_opencl_error(err)));
}
/// Enqueue the kernel for execution and send the mem_refs relating to the
/// results to the next actor. A callback is set to clean up the commmand
/// once the execution is finished. Only called if the results only consist
/// of mem_ref types.
template <class Q = result_types>
detail::enable_if_t<detail::tl_forall<Q, is_ref_type>::value>
enqueue() {
// Errors in this function can not be handled by opencl_err.hpp
// because they require non-standard error handling
CAF_LOG_TRACE("");
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;
auto success = invoke_cl(
clEnqueueNDRangeKernel, parent->queue_.get(), parent->kernel_.get(),
static_cast<cl_uint>(range_.dimensions().size()),
data_or_nullptr(range_.offsets()),
data_or_nullptr(range_.dimensions()),
data_or_nullptr(range_.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 (!success)
return;
auto cb = [](cl_event, cl_int, void* data) {
auto c = reinterpret_cast<command*>(data);
c->deref();
};
if (!invoke_cl(clSetEventCallback, callback_.get(), CL_COMPLETE,
std::move(cb), this))
return;
if (clFlush(parent->queue_.get()) != CL_SUCCESS)
CAF_LOG_ERROR("clFlush: " << CAF_ARG(get_opencl_error(err)));
auto msg = msg_adding_event{callback_}(results_);
promise_.deliver(std::move(msg));
}
private:
template <long I, class T>
void enqueue_read(std::vector<T>&, std::vector<cl_event>& events,
size_t& pos) {
auto p = static_cast<Actor*>(actor_cast<abstract_actor*>(cl_actor_));
events.emplace_back();
auto size = lengths_[pos];
auto buffer_size = sizeof(T) * size;
std::get<I>(results_).resize(size);
auto err = clEnqueueReadBuffer(p->queue_.get(), output_buffers_[pos].get(),
CL_FALSE, 0, buffer_size,
std::get<I>(results_).data(), 1,
events.data(), &events.back());
if (err != CL_SUCCESS) {
this->deref(); // failed to enqueue command
throw std::runtime_error("clEnqueueReadBuffer: " + opencl_error(err));
}
pos += 1;
}
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() {
auto parent = static_cast<Actor*>(actor_cast<abstract_actor*>(cl_actor_));
auto& map_fun = parent->map_results_;
auto msg = map_fun ? apply_args(map_fun, detail::get_indices(results_),
results_)
: message_from_results{}(results_);
promise_.deliver(std::move(msg));
}
// call function F and derefenrence the command on failure
template <class F, class... Us>
bool invoke_cl(F f, Us&&... xs) {
auto err = f(std::forward<Us>(xs)...);
if (err == CL_SUCCESS)
return true;
CAF_LOG_ERROR("error: " << opencl_error(err));
this->deref();
return false;
}
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_;
detail::raw_event_ptr callback_;
std::vector<detail::raw_mem_ptr> input_buffers_;
std::vector<detail::raw_mem_ptr> output_buffers_;
std::vector<detail::raw_mem_ptr> scratch_buffers_;
std::tuple<Ts...> results_;
message msg_; // keeps the argument buffers alive for async copy to device
nd_range range_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_COMMAND_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_DETAIL_COMMAND_HELPER_HPP
#define CAF_OPENCL_DETAIL_COMMAND_HELPER_HPP
#include "caf/detail/type_list.hpp"
#include "caf/opencl/arguments.hpp"
namespace caf {
namespace opencl {
namespace detail {
// 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&...)>;
};
// 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...>;
};
} // namespace detail
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_DETAIL_COMMAND_HELPER_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_CORE_HPP
#define CAF_OPENCL_CORE_HPP
namespace caf {
namespace detail {
} // namespace detail
namespace opencl {
namespace detail {
using namespace caf::detail;
} // namespace detail
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_CORE_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_SMART_PTR_HPP
#define CAF_OPENCL_SMART_PTR_HPP
#include <memory>
#include <algorithm>
#include <type_traits>
#include "caf/intrusive_ptr.hpp"
#include "caf/opencl/global.hpp"
#define CAF_OPENCL_PTR_ALIAS(aliasname, cltype, claddref, clrelease) \
inline void intrusive_ptr_add_ref(cltype ptr) { claddref(ptr); } \
inline void intrusive_ptr_release(cltype ptr) { clrelease(ptr); } \
namespace caf { \
namespace opencl { \
namespace detail { \
using aliasname = intrusive_ptr<std::remove_pointer<cltype>::type>; \
} /* namespace detail */ \
} /* namespace opencl */ \
} // namespace caf
CAF_OPENCL_PTR_ALIAS(raw_mem_ptr, cl_mem, clRetainMemObject, clReleaseMemObject)
CAF_OPENCL_PTR_ALIAS(raw_event_ptr, cl_event, clRetainEvent, clReleaseEvent)
CAF_OPENCL_PTR_ALIAS(raw_kernel_ptr, cl_kernel, clRetainKernel, clReleaseKernel)
CAF_OPENCL_PTR_ALIAS(raw_context_ptr, cl_context,
clRetainContext, clReleaseContext)
CAF_OPENCL_PTR_ALIAS(raw_program_ptr, cl_program,
clRetainProgram, clReleaseProgram)
CAF_OPENCL_PTR_ALIAS(raw_device_ptr, cl_device_id,
clRetainDeviceDummy, clReleaseDeviceDummy)
CAF_OPENCL_PTR_ALIAS(raw_command_queue_ptr, cl_command_queue,
clRetainCommandQueue, clReleaseCommandQueue)
#endif // CAF_OPENCL_SMART_PTR_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_DETAIL_SPAWN_HELPER_HPP
#define CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
#include "caf/opencl/actor_facade.hpp"
namespace caf {
namespace opencl {
namespace detail {
struct tuple_construct { };
template <bool PassConfig, class... Ts>
struct cl_spawn_helper {
using impl = opencl::actor_facade<PassConfig, Ts...>;
using map_in_fun = typename impl::input_mapping;
using map_out_fun = typename impl::output_mapping;
actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::nd_range& range,
Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg),
p, fn, range,
map_in_fun{}, map_out_fun{},
std::forward<Ts>(xs)...));
}
actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::nd_range& range,
map_in_fun map_input, Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg),
p, fn, range,
std::move(map_input),
map_out_fun{},
std::forward<Ts>(xs)...));
}
actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::nd_range& range,
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, range,
std::move(map_input),
std::move(map_output),
std::forward<Ts>(xs)...));
}
};
} // namespace detail
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_DETAIL_SPAWN_HELPER_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_DEVICE_HPP
#define CAF_OPENCL_DEVICE_HPP
#include <vector>
#include "caf/sec.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace caf {
namespace opencl {
class program;
class manager;
template <class T> class mem_ref;
class device;
using device_ptr = intrusive_ptr<device>;
class device : public ref_counted {
public:
friend class program;
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);
detail::raw_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, 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, 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});
}
/// Initialize a new device in a context using a specific device_id
static device_ptr create(const detail::raw_context_ptr& context,
const detail::raw_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
inline unsigned id() const;
/// Returns device info on CL_DEVICE_ADDRESS_BITS
inline cl_uint address_bits() const;
/// Returns device info on CL_DEVICE_ENDIAN_LITTLE
inline cl_bool little_endian() const;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
inline cl_ulong global_mem_cache_size() const;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
inline cl_uint global_mem_cacheline_size() const;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_SIZE
inline cl_ulong global_mem_size() const;
/// Returns device info on CL_DEVICE_HOST_UNIFIED_MEMORY
inline cl_bool host_unified_memory() const;
/// Returns device info on CL_DEVICE_LOCAL_MEM_SIZE
inline cl_ulong local_mem_size() const;
/// Returns device info on CL_DEVICE_LOCAL_MEM_TYPE
inline cl_ulong local_mem_type() const;
/// Returns device info on CL_DEVICE_MAX_CLOCK_FREQUENCY
inline cl_uint max_clock_frequency() const;
/// Returns device info on CL_DEVICE_MAX_COMPUTE_UNITS
inline cl_uint max_compute_units() const;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_ARGS
inline cl_uint max_constant_args() const;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
inline cl_ulong max_constant_buffer_size() const;
/// Returns device info on CL_DEVICE_MAX_MEM_ALLOC_SIZE
inline cl_ulong max_mem_alloc_size() const;
/// Returns device info on CL_DEVICE_MAX_PARAMETER_SIZE
inline size_t max_parameter_size() const;
/// Returns device info on CL_DEVICE_MAX_WORK_GROUP_SIZE
inline size_t max_work_group_size() const;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
inline cl_uint max_work_item_dimensions() const;
/// Returns device info on CL_DEVICE_PROFILING_TIMER_RESOLUTION
inline size_t profiling_timer_resolution() const;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_SIZES
inline const dim_vec& max_work_item_sizes() const;
/// Returns device info on CL_DEVICE_TYPE
inline device_type type() const;
/// Returns device info on CL_DEVICE_EXTENSIONS
inline const std::vector<std::string>& extensions() const;
/// Returns device info on CL_DEVICE_OPENCL_C_VERSION
inline const std::string& opencl_c_version() const;
/// Returns device info on CL_DEVICE_VENDOR
inline const std::string& device_vendor() const;
/// Returns device info on CL_DEVICE_VERSION
inline const std::string& device_version() const;
/// Returns device info on CL_DRIVER_VERSION
inline const std::string& driver_version() const;
/// Returns device info on CL_DEVICE_NAME
inline const std::string& name() const;
private:
device(detail::raw_device_ptr device_id, detail::raw_command_queue_ptr queue,
detail::raw_context_ptr context, unsigned id);
template <class T>
static T info(const detail::raw_device_ptr& device_id, unsigned info_flag) {
T value;
clGetDeviceInfo(device_id.get(), info_flag, sizeof(T), &value, nullptr);
return value;
}
static std::string info_string(const detail::raw_device_ptr& device_id,
unsigned info_flag);
detail::raw_device_ptr device_id_;
detail::raw_command_queue_ptr queue_;
detail::raw_context_ptr context_;
unsigned id_;
bool profiling_enabled_; // CL_DEVICE_QUEUE_PROPERTIES
bool out_of_order_execution_; // CL_DEVICE_QUEUE_PROPERTIES
cl_uint address_bits_; // CL_DEVICE_ADDRESS_BITS
cl_bool little_endian_; // CL_DEVICE_ENDIAN_LITTLE
cl_ulong global_mem_cache_size_; // CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
cl_uint global_mem_cacheline_size_; // CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
cl_ulong global_mem_size_; // CL_DEVICE_GLOBAL_MEM_SIZE
cl_bool host_unified_memory_; // CL_DEVICE_HOST_UNIFIED_MEMORY
cl_ulong local_mem_size_; // CL_DEVICE_LOCAL_MEM_SIZE
cl_uint local_mem_type_; // CL_DEVICE_LOCAL_MEM_TYPE
cl_uint max_clock_frequency_; // CL_DEVICE_MAX_CLOCK_FREQUENCY
cl_uint max_compute_units_; // CL_DEVICE_MAX_COMPUTE_UNITS
cl_uint max_constant_args_; // CL_DEVICE_MAX_CONSTANT_ARGS
cl_ulong max_constant_buffer_size_; // CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
cl_ulong max_mem_alloc_size_; // CL_DEVICE_MAX_MEM_ALLOC_SIZE
size_t max_parameter_size_; // CL_DEVICE_MAX_PARAMETER_SIZE
size_t max_work_group_size_; // CL_DEVICE_MAX_WORK_GROUP_SIZE
cl_uint max_work_item_dimensions_; // CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
size_t profiling_timer_resolution_; // CL_DEVICE_PROFILING_TIMER_RESOLUTION
dim_vec max_work_item_sizes_; // CL_DEVICE_MAX_WORK_ITEM_SIZES
device_type device_type_; // CL_DEVICE_TYPE
std::vector<std::string> extensions_; // CL_DEVICE_EXTENSIONS
std::string opencl_c_version_; // CL_DEVICE_OPENCL_C_VERSION
std::string device_vendor_; // CL_DEVICE_VENDOR
std::string device_version_; // CL_DEVICE_VERSION
std::string driver_version_; // CL_DRIVER_VERSION
std::string name_; // CL_DEVICE_NAME
};
/******************************************************************************\
* implementation of inline member functions *
\******************************************************************************/
inline unsigned device::id() const {
return id_;
}
inline cl_uint device::address_bits() const {
return address_bits_;
}
inline cl_bool device::little_endian() const {
return little_endian_;
}
inline cl_ulong device::global_mem_cache_size() const {
return global_mem_cache_size_;
}
inline cl_uint device::global_mem_cacheline_size() const {
return global_mem_cacheline_size_;
}
inline cl_ulong device::global_mem_size() const {
return global_mem_size_;
}
inline cl_bool device::host_unified_memory() const {
return host_unified_memory_;
}
inline cl_ulong device::local_mem_size() const {
return local_mem_size_;
}
inline cl_ulong device::local_mem_type() const {
return local_mem_size_;
}
inline cl_uint device::max_clock_frequency() const {
return max_clock_frequency_;
}
inline cl_uint device::max_compute_units() const {
return max_compute_units_;
}
inline cl_uint device::max_constant_args() const {
return max_constant_args_;
}
inline cl_ulong device::max_constant_buffer_size() const {
return max_constant_buffer_size_;
}
inline cl_ulong device::max_mem_alloc_size() const {
return max_mem_alloc_size_;
}
inline size_t device::max_parameter_size() const {
return max_parameter_size_;
}
inline size_t device::max_work_group_size() const {
return max_work_group_size_;
}
inline cl_uint device::max_work_item_dimensions() const {
return max_work_item_dimensions_;
}
inline size_t device::profiling_timer_resolution() const {
return profiling_timer_resolution_;
}
inline const dim_vec& device::max_work_item_sizes() const {
return max_work_item_sizes_;
}
inline device_type device::type() const {
return device_type_;
}
inline const std::vector<std::string>& device::extensions() const {
return extensions_;
}
inline const std::string& device::opencl_c_version() const {
return opencl_c_version_;
}
inline const std::string& device::device_vendor() const {
return device_vendor_;
}
inline const std::string& device::device_version() const {
return device_version_;
}
inline const std::string& device::driver_version() const {
return driver_version_;
}
inline const std::string& device::name() const {
return name_;
}
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_DEVICE_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_GLOBAL_HPP
#define CAF_OPENCL_GLOBAL_HPP
#include <string>
#include <iostream>
#include "caf/config.hpp"
#include "caf/detail/limited_vector.hpp"
#if defined(CAF_MACOS) || defined(CAF_IOS)
# include <OpenCL/opencl.h>
#else
# include <CL/opencl.h>
#endif
// needed for OpenCL 1.0 compatibility (works around missing clReleaseDevice)
extern "C" {
cl_int clReleaseDeviceDummy(cl_device_id);
cl_int clRetainDeviceDummy(cl_device_id);
} // extern "C"
namespace caf {
namespace opencl {
enum device_type {
def = CL_DEVICE_TYPE_DEFAULT,
cpu = CL_DEVICE_TYPE_CPU,
gpu = CL_DEVICE_TYPE_GPU,
accelerator = CL_DEVICE_TYPE_ACCELERATOR,
custom = CL_DEVICE_TYPE_CUSTOM,
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);
device_type device_type_from_ulong(cl_ulong dev);
/// A vector of up to three elements used for OpenCL dimensions.
using dim_vec = detail::limited_vector<size_t, 3>;
std::string opencl_error(cl_int err);
std::string event_status(cl_event event);
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_GLOBAL_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_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/actor_facade.hpp"
#include "caf/opencl/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.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 detail::raw_command_queue_ptr command_queue(uint32_t id);
manager(const manager&) = delete;
manager& operator=(const manager&) = delete;
/// Get the device with id, which is assigned sequientally.
optional<device_ptr> find_device(size_t dev_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>
optional<device_ptr> find_device_if(UnaryPredicate p) const {
for (auto& pl : platforms_) {
for (auto& dev : pl->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,
detail::type_list<>);
// 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
/// from a given @p kernel_source.
/// @returns A program object.
program_ptr create_program(const char* kernel_source,
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
/// from a given @p kernel_source.
/// @returns A program object.
program_ptr create_program(const char* kernel_source,
const char* options, const device_ptr 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>
detail::enable_if_t<opencl::is_opencl_arg<T>::value, actor>
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::nd_range& range, T&& x, Ts&&... xs) {
detail::cl_spawn_helper<false, T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, range,
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>
detail::enable_if_t<opencl::is_opencl_arg<T>::value, actor>
spawn(const char* source, const char* fname,
const opencl::nd_range& range, T&& x, Ts&&... xs) {
detail::cl_spawn_helper<false, T, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, range,
std::forward<T>(x), 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 Fun, class... Ts>
actor spawn(const opencl::program_ptr prog, const char* fname,
const opencl::nd_range& range,
std::function<optional<message> (message&)> map_args,
Fun map_result, Ts&&... xs) {
detail::cl_spawn_helper<false, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, range,
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::nd_range& range,
std::function<optional<message> (message&)> map_args,
Fun map_result, Ts&&... xs) {
detail::cl_spawn_helper<false, Ts...> f;
return f(actor_config{system_.dummy_execution_unit()},
create_program(source), fname, range,
std::move(map_args), std::move(map_result),
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>
detail::enable_if_t<opencl::is_opencl_arg<T>::value, actor>
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::nd_range& range,
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, range,
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>
detail::enable_if_t<opencl::is_opencl_arg<T>::value, actor>
spawn(const char* source, const char* fname,
const opencl::nd_range& range,
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, range,
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>
detail::enable_if_t<!opencl::is_opencl_arg<Fun>::value, actor>
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::nd_range& range,
std::function<optional<message> (nd_range&, 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, range,
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>
detail::enable_if_t<!opencl::is_opencl_arg<Fun>::value, actor>
spawn(const char* source, const char* fname,
const opencl::nd_range& range,
std::function<optional<message> (nd_range&, 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, range,
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>
detail::enable_if_t<opencl::is_opencl_arg<T>::value, actor>
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::nd_range& range,
std::function<optional<message> (nd_range&, 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, range,
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>
detail::enable_if_t<opencl::is_opencl_arg<T>::value, actor>
spawn(const char* source, const char* fname,
const opencl::nd_range& range,
std::function<optional<message> (nd_range&, 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, range,
std::move(map_args), std::forward<T>(x), std::forward<Ts>(xs)...);
}
protected:
manager(actor_system& sys);
~manager() override;
private:
actor_system& system_;
std::vector<platform_ptr> platforms_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_MANAGER_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_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/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace caf {
namespace opencl {
/// Updates the reference types in a message with a given event.
struct msg_adding_event {
msg_adding_event(detail::raw_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);
}
detail::raw_event_ptr event_;
};
// Tag to separate mem_refs from other types in messages.
struct ref_tag {};
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 : ref_tag {
public:
using value_type = T;
friend struct msg_adding_event;
template <bool PassConfig, class... Ts>
friend class actor_facade;
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, 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 detail::raw_mem_ptr& get() const {
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, detail::raw_command_queue_ptr queue,
detail::raw_mem_ptr memory, cl_mem_flags access,
detail::raw_event_ptr event)
: num_elements_{num_elements},
access_{access},
queue_{queue},
event_{event},
memory_{memory} {
// nop
}
mem_ref(size_t num_elements, detail::raw_command_queue_ptr queue,
cl_mem memory, cl_mem_flags access, detail::raw_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(detail::raw_event_ptr e) {
event_ = std::move(e);
}
inline detail::raw_event_ptr event() {
return event_;
}
inline cl_event take_event() {
return event_.release();
}
size_t num_elements_;
cl_mem_flags access_;
detail::raw_command_queue_ptr queue_;
detail::raw_event_ptr event_;
detail::raw_mem_ptr memory_;
};
} // namespace opencl
template <class T>
struct allowed_unsafe_message_type<opencl::mem_ref<T>> : std::true_type {};
} // 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_ND_RANGE_HPP
#define CAF_OPENCL_ND_RANGE_HPP
#include "caf/opencl/global.hpp"
namespace caf {
namespace opencl {
class nd_range {
public:
nd_range(const opencl::dim_vec& dimensions,
const opencl::dim_vec& offsets = {},
const opencl::dim_vec& local_dimensions = {})
: dims_{dimensions},
offset_{offsets},
local_dims_{local_dimensions} {
// nop
}
nd_range(opencl::dim_vec&& dimensions,
opencl::dim_vec&& offsets = {},
opencl::dim_vec&& local_dimensions = {})
: dims_{std::move(dimensions)},
offset_{std::move(offsets)},
local_dims_{std::move(local_dimensions)} {
// nop
}
nd_range(const nd_range&) = default;
nd_range(nd_range&&) = default;
nd_range& operator=(const nd_range&) = default;
nd_range& operator=(nd_range&&) = default;
const opencl::dim_vec& dimensions() const {
return dims_;
}
const opencl::dim_vec& offsets() const {
return offset_;
}
const opencl::dim_vec& local_dimensions() const {
return local_dims_;
}
private:
opencl::dim_vec dims_;
opencl::dim_vec offset_;
opencl::dim_vec local_dims_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_ND_RANGE_HPP
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* Raphael Hiesgen <raphael.hiesgen (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_OPENCL_ERR_HPP
#define CAF_OPENCL_OPENCL_ERR_HPP
#include "caf/opencl/global.hpp"
#include "caf/logger.hpp"
#define CAF_CLF(funname) #funname , funname
namespace caf {
namespace opencl {
void throwcl(const char* fname, cl_int err);
void pfn_notify(const char* errinfo, const void*, size_t, void*);
// call convention for simply calling a function
template <class F, class... Ts>
void v1callcl(const char* fname, F f, Ts&&... vs) {
throwcl(fname, f(std::forward<Ts>(vs)...));
}
// call convention for simply calling a function, not using the last argument
template <class F, class... Ts>
void v2callcl(const char* fname, F f, Ts&&... vs) {
throwcl(fname, f(std::forward<Ts>(vs)..., nullptr));
}
// call convention with `result` argument at the end returning `err`, not
// using the second last argument (set to nullptr) nor the one before (set to 0)
template <class R, class F, class... Ts>
R v1get(const char* fname, F f, Ts&&... vs) {
R result;
throwcl(fname, f(std::forward<Ts>(vs)..., cl_uint{0}, nullptr, &result));
return result;
}
// call convention with `err` argument at the end returning `result`
template <class F, class... Ts>
auto v2get(const char* fname, F f, Ts&&... vs)
-> decltype(f(std::forward<Ts>(vs)..., nullptr)) {
cl_int err;
auto result = f(std::forward<Ts>(vs)..., &err);
throwcl(fname, err);
return result;
}
// call convention with `result` argument at second last position (preceeded by
// its size) followed by an ingored void* argument (nullptr) returning `err`
template <class R, class F, class... Ts>
R v3get(const char* fname, F f, Ts&&... vs) {
R result;
throwcl(fname, f(std::forward<Ts>(vs)..., sizeof(R), &result, nullptr));
return result;
}
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_OPENCL_ERR_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_PLATFORM_HPP
#define CAF_OPENCL_PLATFORM_HPP
#include "caf/ref_counted.hpp"
#include "caf/opencl/device.hpp"
namespace caf {
namespace opencl {
class platform;
using platform_ptr = intrusive_ptr<platform>;
class platform : public ref_counted {
public:
friend class program;
template <class T, class... Ts>
friend intrusive_ptr<T> caf::make_counted(Ts&&...);
inline const std::vector<device_ptr>& devices() const;
inline const std::string& name() const;
inline const std::string& vendor() const;
inline const std::string& version() const;
static platform_ptr create(cl_platform_id platform_id, unsigned start_id);
private:
platform(cl_platform_id platform_id, detail::raw_context_ptr context,
std::string name, std::string vendor, std::string version,
std::vector<device_ptr> devices);
~platform();
static std::string platform_info(cl_platform_id platform_id,
unsigned info_flag);
cl_platform_id platform_id_;
detail::raw_context_ptr context_;
std::string name_;
std::string vendor_;
std::string version_;
std::vector<device_ptr> devices_;
};
/******************************************************************************\
* implementation of inline member functions *
\******************************************************************************/
inline const std::vector<device_ptr>& platform::devices() const {
return devices_;
}
inline const std::string& platform::name() const {
return name_;
}
inline const std::string& platform::vendor() const {
return vendor_;
}
inline const std::string& platform::version() const {
return version_;
}
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_PLATFORM_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_PROGRAM_HPP
#define CAF_OPENCL_PROGRAM_HPP
#include <map>
#include <memory>
#include "caf/ref_counted.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace caf {
namespace opencl {
class program;
using program_ptr = intrusive_ptr<program>;
/// @brief A wrapper for OpenCL's cl_program.
class program : public ref_counted {
public:
friend class manager;
template <bool PassConfig, class... Ts>
friend class actor_facade;
template <class T, class... Ts>
friend intrusive_ptr<T> caf::make_counted(Ts&&...);
private:
program(detail::raw_context_ptr context, detail::raw_command_queue_ptr queue,
detail::raw_program_ptr prog,
std::map<std::string, detail::raw_kernel_ptr> available_kernels);
~program();
detail::raw_context_ptr context_;
detail::raw_program_ptr program_;
detail::raw_command_queue_ptr queue_;
std::map<std::string, detail::raw_kernel_ptr> available_kernels_;
};
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_PROGRAM_HPP
cmake_minimum_required(VERSION 2.8)
project(caf_examples_opencl CXX)
if(OpenCL_LIBRARIES AND NOT CAF_NO_EXAMPLES)
add_custom_target(opencl_examples)
include_directories(${LIBCAF_INCLUDE_DIRS})
if(${CMAKE_SYSTEM_NAME} MATCHES "Window")
set(WSLIB -lws2_32)
else()
set(WSLIB)
endif()
macro(add name folder)
add_executable(${name} ${folder}/${name}.cpp ${ARGN})
target_link_libraries(${name}
${LD_FLAGS}
${CAF_LIBRARIES}
${PTHREAD_LIBRARIES}
${WSLIB}
${OpenCL_LIBRARIES})
add_dependencies(${name} all_examples)
endmacro()
add(proper_matrix .)
add(simple_matrix .)
add(scan .)
endif()
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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 <iomanip>
#include <numeric>
#include <cassert>
#include <iostream>
#include "caf/all.hpp"
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/all.hpp"
using namespace std;
using namespace caf;
using namespace caf::opencl;
using caf::detail::limited_vector;
namespace {
using fvec = vector<float>;
constexpr size_t matrix_size = 8;
constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2
// last parameter is, by convention, the output parameter
constexpr const char* kernel_source = R"__(
kernel void matrix_mult(global const float* matrix1,
global const float* matrix2,
global float* output) {
// we only use square matrices, hence: width == height
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);
float result = 0;
for (size_t idx = 0; idx < size; ++idx)
result += matrix1[idx + y * size] * matrix2[x + idx * size];
output[x+y*size] = result;
}
)__";
} // namespace <anonymous>
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 Inspector>
friend typename Inspector::result_type inspect(Inspector& f,
square_matrix& m) {
return f(meta::type_name("square_matrix"), m.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) { }
explicit square_matrix(fvec d) : data_(move(d)) {
assert(data_.size() == num_elements);
}
inline float& operator()(size_t column, size_t row) {
return data_[column + row * Size];
}
inline const float& operator()(size_t column, size_t row) const {
return data_[column + row * Size];
}
inline void iota_fill() {
iota(data_.begin(), data_.end(), 0);
}
using const_iterator = typename fvec::const_iterator;
const_iterator begin() const { return data_.begin(); }
const_iterator end() const { return data_.end(); }
fvec& data() { return data_; }
const fvec& data() const { return data_; }
private:
fvec data_;
};
template<size_t Size>
string to_string(const square_matrix<Size>& m) {
ostringstream oss;
oss.fill(' ');
for (size_t row = 0; row < Size; ++row) {
for (size_t column = 0; column < Size; ++column)
oss << fixed << setprecision(2) << setw(9) << m(column, row);
oss << '\n';
}
return oss.str();
}
// to annouce the square_matrix to libcaf, we have to define these operators
template<size_t Size>
inline bool operator==(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) {
return equal(lhs.begin(), lhs.end(), rhs.begin());
}
template<size_t Size>
inline bool operator!=(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) {
return !(lhs == rhs);
}
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;
m1.iota_fill();
auto m2 = m1;
// print "source" matrix
cout << "calculating square of matrix:" << endl
<< to_string(m1) << endl;
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()));
});
};
auto box_res = [] (fvec& result) -> message {
return make_message(matrix_type{move(result)});
};
// spawn an opencl actor
// 1st arg: source code of one or more opencl kernels
// 2nd arg: name of the kernel to use
// 3rd arg: the config specifies how many dimensions the kernel uses and how
// many work items are created, creates matrix_size * matrix_size
// global work items in this case
// 4th arg: the opencl function operates on vectors, this function converts
// a tuple of two matrices to a tuple of vectors; note that this
// function returns an option (an empty results causes the actor to
// ignore the message)
// 5th arg: converts the ouptut vector back to a matrix that is then
// used as response message
// from 6 : a description of the kernel signature using in/out/in_out classes
// 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,
nd_range{dim_vec{matrix_size, matrix_size}},
unbox_args, box_res,
in<float>{}, in<float>{}, out<float>{});
// send both matrices to the actor and
// wait for results in form of a matrix_type
self->request(worker, chrono::seconds(5), move(m1), move(m2)).then(
[](const matrix_type& result) {
cout << "result:" << endl << to_string(result);
}
);
}
int main() {
// matrix_type ist not a simple type,
// it must be annouced to libcaf
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;
}
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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>
template <class T, class E = caf::detail::enable_if_t<is_integral<T>::value>>
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(problem_size);
random_device rd;
mt19937 gen(rd());
uniform_int_distribution<uval> val_gen(0, 1023);
std::generate(begin(values), end(values), [&]() { return val_gen(gen); });
// ---- find device ----
auto& mngr = system.opencl_manager();
//
string prefix = "GeForce";
auto opt = mngr.find_device_if([&](const device_ptr dev) {
auto& name = dev->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.find_device();
}
if (!opt) {
cerr << "Not OpenCL device available." << endl;
return 0;
} else {
cerr << "Found device '" << (*opt)->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->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 nd_range{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 = nd_range{dim_vec{half_block}, {}, dim_vec{half_block}};
// ---- scan actors ----
auto phase1 = mngr.spawn(
prog, kernel_name_1, ndr,
[nd_conf](nd_range& range, message& msg) -> optional<message> {
return msg.apply([&](uvec& vec) {
auto size = vec.size();
range = 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](nd_range& range, message& msg) -> optional<message> {
return msg.apply([&](uref& data, uref& incs) {
auto size = incs.size();
range = 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](nd_range& range, message& msg) -> optional<message> {
return msg.apply([&](uref& data, uref& incs) {
auto size = incs.size();
range = 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;
}
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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 <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 fvec = std::vector<float>;
constexpr size_t matrix_size = 8;
constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2
// last parameter is, by convention, the output parameter
constexpr const char* kernel_source = R"__(
kernel void matrix_mult(global const float* matrix1,
global const float* matrix2,
global float* output) {
// we only use square matrices, hence: width == height
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);
float result = 0;
for (size_t idx = 0; idx < size; ++idx)
result += matrix1[idx + y * size] * matrix2[x + idx * size];
output[x+y*size] = result;
}
)__";
} // namespace <anonymous>
void print_as_matrix(const fvec& matrix) {
for (size_t column = 0; column < matrix_size; ++column) {
for (size_t row = 0; row < matrix_size; ++row) {
cout << fixed << setprecision(2) << setw(9)
<< matrix[row + column * matrix_size];
}
cout << endl;
}
}
void multiplier(event_based_actor* self) {
// the opencl actor only understands vectors
// so these vectors represent the matrices
fvec m1(matrix_size * matrix_size);
fvec m2(matrix_size * matrix_size);
// fill each with ascending values
iota(m1.begin(), m1.end(), 0);
iota(m2.begin(), m2.end(), 0);
// print "source" matrix
cout << "calculating square of matrix:" << endl;
print_as_matrix(m1);
cout << endl;
// spawn an opencl actor
// 1st arg: source code of one or more kernels
// 2nd arg: name of the kernel to use
// 3rd arg: a spawn configuration that includes:
// - the global dimension arguments for opencl's enqueue
// creates matrix_size * matrix_size global work items
// - offsets for global dimensions (optional)
// - local dimensions (optional)
// 4th to Nth arg: the kernel signature described by in/out/in_out classes
// that contain the argument type in their template. 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(
kernel_source, kernel_name,
nd_range{dim_vec{matrix_size, matrix_size}},
in<float>{}, in<float>{}, out<float>{}
);
// send both matrices to the actor and wait for a result
self->request(worker, chrono::seconds(5), move(m1), move(m2)).then(
[](const fvec& result) {
cout << "result: " << endl;
print_as_matrix(result);
}
);
}
int main() {
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;
}
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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 <iostream>
#include <utility>
#include "caf/logger.hpp"
#include "caf/ref_counted.hpp"
#include "caf/string_algorithms.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/opencl_err.hpp"
using namespace std;
namespace caf {
namespace opencl {
device_ptr device::create(const detail::raw_context_ptr& context,
const detail::raw_device_ptr& device_id,
unsigned 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 = false; // (supported & CL_QUEUE_PROFILING_ENABLE) != 0u;
bool out_of_order = (supported & CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE) != 0u;
unsigned properties = profiling ? CL_QUEUE_PROFILING_ENABLE : 0;
if (out_of_order)
properties |= CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE;
// create the command queue
detail::raw_command_queue_ptr command_queue{
v2get(CAF_CLF(clCreateCommandQueue),
context.get(), device_id.get(),
properties),
false
};
// create the device
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
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->global_mem_cache_size_ =
info<cl_ulong>(device_id, CL_DEVICE_GLOBAL_MEM_CACHE_SIZE);
dev->global_mem_cacheline_size_ =
info<cl_uint>(device_id, CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE);
dev->global_mem_size_ =
info<cl_ulong >(device_id,CL_DEVICE_GLOBAL_MEM_SIZE);
dev->host_unified_memory_ =
info<cl_bool>(device_id, CL_DEVICE_HOST_UNIFIED_MEMORY);
dev->local_mem_size_ =
info<cl_ulong>(device_id, CL_DEVICE_LOCAL_MEM_SIZE);
dev->local_mem_type_ =
info<cl_uint>(device_id, CL_DEVICE_LOCAL_MEM_TYPE);
dev->max_clock_frequency_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_CLOCK_FREQUENCY);
dev->max_compute_units_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_COMPUTE_UNITS);
dev->max_constant_args_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_CONSTANT_ARGS);
dev->max_constant_buffer_size_ =
info<cl_ulong>(device_id, CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE);
dev->max_mem_alloc_size_ =
info<cl_ulong>(device_id, CL_DEVICE_MAX_MEM_ALLOC_SIZE);
dev->max_parameter_size_ =
info<size_t>(device_id, CL_DEVICE_MAX_PARAMETER_SIZE);
dev->max_work_group_size_ =
info<size_t>(device_id, CL_DEVICE_MAX_WORK_GROUP_SIZE);
dev->max_work_item_dimensions_ =
info<cl_uint>(device_id, CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS);
dev->profiling_timer_resolution_ =
info<size_t>(device_id, CL_DEVICE_PROFILING_TIMER_RESOLUTION);
dev->max_work_item_sizes_.resize(dev->max_work_item_dimensions_);
clGetDeviceInfo(device_id.get(), CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(size_t) * dev->max_work_item_dimensions_,
dev->max_work_item_sizes_.data(), nullptr);
dev->device_type_ =
device_type_from_ulong(info<cl_ulong>(device_id, CL_DEVICE_TYPE));
string extensions = info_string(device_id, CL_DEVICE_EXTENSIONS);
split(dev->extensions_, extensions, " ", false);
dev->opencl_c_version_ = info_string(device_id, CL_DEVICE_EXTENSIONS);
dev->device_vendor_ = info_string(device_id, CL_DEVICE_VENDOR);
dev->device_version_ = info_string(device_id, CL_DEVICE_VERSION);
dev->name_ = info_string(device_id, CL_DEVICE_NAME);
return dev;
}
void device::synchronize() {
clFinish(queue_.get());
}
string device::info_string(const detail::raw_device_ptr& device_id,
unsigned info_flag) {
size_t size;
clGetDeviceInfo(device_id.get(), info_flag, 0, nullptr, &size);
vector<char> buffer(size);
clGetDeviceInfo(device_id.get(), info_flag, sizeof(char) * size, buffer.data(),
nullptr);
return string(buffer.data());
}
device::device(detail::raw_device_ptr device_id,
detail::raw_command_queue_ptr queue,
detail::raw_context_ptr context,
unsigned id)
: device_id_(std::move(device_id)),
queue_(std::move(queue)),
context_(std::move(context)),
id_(id) {
// nop
}
device::~device() {
// nop
}
} // namespace opencl
} // namespace caf
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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 <string>
#include <sstream>
#include "caf/opencl/global.hpp"
cl_int clReleaseDeviceDummy(cl_device_id) {
return 0;
}
cl_int clRetainDeviceDummy(cl_device_id) {
return 0;
}
namespace caf {
namespace opencl {
std::ostream& operator<<(std::ostream& os, device_type dev) {
switch(dev) {
case def : os << "default"; break;
case cpu : os << "CPU"; break;
case gpu : os << "GPU"; break;
case accelerator : os << "accelerator"; break;
case custom : os << "custom"; break;
case all : os << "all"; break;
default : os.setstate(std::ios_base::failbit);
}
return os;
}
device_type device_type_from_ulong(cl_ulong dev) {
switch(dev) {
case CL_DEVICE_TYPE_CPU : return cpu;
case CL_DEVICE_TYPE_GPU : return gpu;
case CL_DEVICE_TYPE_ACCELERATOR : return accelerator;
case CL_DEVICE_TYPE_CUSTOM : return custom;
case CL_DEVICE_TYPE_ALL : return all;
default : return def;
}
}
std::string opencl_error(cl_int err) {
switch (err) {
case CL_SUCCESS:
return "CL_SUCCESS";
case CL_DEVICE_NOT_FOUND:
return "CL_DEVICE_NOT_FOUND";
case CL_DEVICE_NOT_AVAILABLE:
return "CL_DEVICE_NOT_AVAILABLE";
case CL_COMPILER_NOT_AVAILABLE:
return "CL_COMPILER_NOT_AVAILABLE";
case CL_MEM_OBJECT_ALLOCATION_FAILURE:
return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
case CL_OUT_OF_RESOURCES:
return "CL_OUT_OF_RESOURCES";
case CL_OUT_OF_HOST_MEMORY:
return "CL_OUT_OF_HOST_MEMORY";
case CL_PROFILING_INFO_NOT_AVAILABLE:
return "CL_PROFILING_INFO_NOT_AVAILABLE";
case CL_MEM_COPY_OVERLAP:
return "CL_MEM_COPY_OVERLAP";
case CL_IMAGE_FORMAT_MISMATCH:
return "CL_IMAGE_FORMAT_MISMATCH";
case CL_IMAGE_FORMAT_NOT_SUPPORTED:
return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
case CL_BUILD_PROGRAM_FAILURE:
return "CL_BUILD_PROGRAM_FAILURE";
case CL_MAP_FAILURE:
return "CL_MAP_FAILURE";
case CL_INVALID_VALUE:
return "CL_INVALID_VALUE";
case CL_INVALID_DEVICE_TYPE:
return "CL_INVALID_DEVICE_TYPE";
case CL_INVALID_PLATFORM:
return "CL_INVALID_PLATFORM";
case CL_INVALID_DEVICE:
return "CL_INVALID_DEVICE";
case CL_INVALID_CONTEXT:
return "CL_INVALID_CONTEXT";
case CL_INVALID_QUEUE_PROPERTIES:
return "CL_INVALID_QUEUE_PROPERTIES";
case CL_INVALID_COMMAND_QUEUE:
return "CL_INVALID_COMMAND_QUEUE";
case CL_INVALID_HOST_PTR:
return "CL_INVALID_HOST_PTR";
case CL_INVALID_MEM_OBJECT:
return "CL_INVALID_MEM_OBJECT";
case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
case CL_INVALID_IMAGE_SIZE:
return "CL_INVALID_IMAGE_SIZE";
case CL_INVALID_SAMPLER:
return "CL_INVALID_SAMPLER";
case CL_INVALID_BINARY:
return "CL_INVALID_BINARY";
case CL_INVALID_BUILD_OPTIONS:
return "CL_INVALID_BUILD_OPTIONS";
case CL_INVALID_PROGRAM:
return "CL_INVALID_PROGRAM";
case CL_INVALID_PROGRAM_EXECUTABLE:
return "CL_INVALID_PROGRAM_EXECUTABLE";
case CL_INVALID_KERNEL_NAME:
return "CL_INVALID_KERNEL_NAME";
case CL_INVALID_KERNEL_DEFINITION:
return "CL_INVALID_KERNEL_DEFINITION";
case CL_INVALID_KERNEL:
return "CL_INVALID_KERNEL";
case CL_INVALID_ARG_INDEX:
return "CL_INVALID_ARG_INDEX";
case CL_INVALID_ARG_VALUE:
return "CL_INVALID_ARG_VALUE";
case CL_INVALID_ARG_SIZE:
return "CL_INVALID_ARG_SIZE";
case CL_INVALID_KERNEL_ARGS:
return "CL_INVALID_KERNEL_ARGS";
case CL_INVALID_WORK_DIMENSION:
return "CL_INVALID_WORK_DIMENSION";
case CL_INVALID_WORK_GROUP_SIZE:
return "CL_INVALID_WORK_GROUP_SIZE";
case CL_INVALID_WORK_ITEM_SIZE:
return "CL_INVALID_WORK_ITEM_SIZE";
case CL_INVALID_GLOBAL_OFFSET:
return "CL_INVALID_GLOBAL_OFFSET";
case CL_INVALID_EVENT_WAIT_LIST:
return "CL_INVALID_EVENT_WAIT_LIST";
case CL_INVALID_EVENT:
return "CL_INVALID_EVENT";
case CL_INVALID_OPERATION:
return "CL_INVALID_OPERATION";
case CL_INVALID_GL_OBJECT:
return "CL_INVALID_GL_OBJECT";
case CL_INVALID_BUFFER_SIZE:
return "CL_INVALID_BUFFER_SIZE";
case CL_INVALID_MIP_LEVEL:
return "CL_INVALID_MIP_LEVEL";
case CL_INVALID_GLOBAL_WORK_SIZE:
return "CL_INVALID_GLOBAL_WORK_SIZE";
// error codes used by extensions
// see: http://streamcomputing.eu/blog/2013-04-28/opencl-1-2-error-codes/
case -1000:
return "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR";
case -1001:
return "No valid ICDs found";
case -1002:
return "CL_INVALID_D3D10_DEVICE_KHR";
case -1003:
return "CL_INVALID_D3D10_RESOURCE_KHR";
case -1004:
return "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR";
case -1005:
return "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR";
default:
return "UNKNOWN_ERROR: " + std::to_string(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 caf
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* Raphael Hiesgen <raphael.hiesgen (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include <fstream>
#include "caf/detail/type_list.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/manager.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
using namespace std;
namespace caf {
namespace opencl {
optional<device_ptr> manager::find_device(size_t dev_id) const {
if (platforms_.empty())
return none;
size_t to = 0;
for (auto& pl : platforms_) {
auto from = to;
to += pl->devices().size();
if (dev_id >= from && dev_id < to)
return pl->devices()[dev_id - from];
}
return none;
}
void manager::init(actor_system_config&) {
// get number of available platforms
auto num_platforms = v1get<cl_uint>(CAF_CLF(clGetPlatformIDs));
// get platform ids
std::vector<cl_platform_id> platform_ids(num_platforms);
v2callcl(CAF_CLF(clGetPlatformIDs), num_platforms, platform_ids.data());
if (platform_ids.empty())
throw std::runtime_error("no OpenCL platform found");
// initialize platforms (device discovery)
unsigned current_device_id = 0;
for (auto& pl_id : platform_ids) {
platforms_.push_back(platform::create(pl_id, current_device_id));
current_device_id +=
static_cast<unsigned>(platforms_.back()->devices().size());
}
}
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_ptr manager::create_program_from_file(const char* path,
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 = find_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_ptr manager::create_program_from_file(const char* path,
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
size_t kernel_source_length = strlen(kernel_source);
detail::raw_program_ptr pptr;
pptr.reset(v2get(CAF_CLF(clCreateProgramWithSource), dev->context_.get(),
1u, &kernel_source, &kernel_source_length),
false);
// build programm from program object
auto dev_tmp = dev->device_id_.get();
auto err = clBuildProgram(pptr.get(), 1, &dev_tmp, options, nullptr, nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clBuildProgram: " << opencl_error(err);
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,
0, 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(char) * buildlog_buffer_size,
buffer.data(), nullptr);
ostringstream ss;
ss << "############## Build log ##############"
<< 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()));
#endif
oss << endl << ss.str();
}
throw runtime_error(oss.str());
}
cl_uint number_of_kernels = 0;
clCreateKernelsInProgram(pptr.get(), 0u, nullptr, &number_of_kernels);
map<string, detail::raw_kernel_ptr> available_kernels;
if (number_of_kernels > 0) {
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: " << opencl_error(err);
throw runtime_error(oss.str());
}
for (cl_uint i = 0; i < number_of_kernels; ++i) {
size_t len;
clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, 0, nullptr, &len);
vector<char> name(len);
err = clGetKernelInfo(kernels[i], CL_KERNEL_FUNCTION_NAME, len,
reinterpret_cast<void*>(name.data()), nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetKernelInfo (CL_KERNEL_FUNCTION_NAME): "
<< opencl_error(err);
throw runtime_error(oss.str());
}
detail::raw_kernel_ptr kernel;
kernel.reset(move(kernels[i]));
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 make_counted<program>(dev->context_, dev->queue_, pptr,
move(available_kernels));
}
manager::manager(actor_system& sys) : system_(sys) {
// nop
}
manager::~manager() {
// nop
}
} // namespace opencl
} // namespace caf
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* Raphael Hiesgen <raphael.hiesgen (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENSE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include "caf/opencl/opencl_err.hpp"
namespace caf {
namespace opencl {
void throwcl(const char* fname, cl_int err) {
if (err != CL_SUCCESS) {
std::string errstr = fname;
errstr += ": ";
errstr += opencl_error(err);
throw std::runtime_error(std::move(errstr));
}
}
void pfn_notify(const char* errinfo, const void*, size_t, void*) {
CAF_LOG_ERROR("\n##### Error message via pfn_notify #####\n"
<< errinfo <<
"\n########################################");
static_cast<void>(errinfo); // remove warning
}
} // namespace opencl
} // namespace caf
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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 <utility>
#include <vector>
#include <iostream>
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
using namespace std;
namespace caf {
namespace opencl {
platform_ptr platform::create(cl_platform_id platform_id,
unsigned start_id) {
vector<unsigned> device_types = {CL_DEVICE_TYPE_GPU,
CL_DEVICE_TYPE_ACCELERATOR,
CL_DEVICE_TYPE_CPU};
vector<cl_device_id> ids;
for (cl_device_type device_type : device_types) {
auto known = ids.size();
cl_uint discoverd;
auto err = clGetDeviceIDs(platform_id, device_type, 0, nullptr, &discoverd);
if (err == CL_DEVICE_NOT_FOUND) {
continue; // no devices of the type found
} else if (err != CL_SUCCESS) {
throwcl("clGetDeviceIDs", err);
}
ids.resize(known + discoverd);
v2callcl(CAF_CLF(clGetDeviceIDs), platform_id, device_type,
discoverd, (ids.data() + known));
}
vector<detail::raw_device_ptr> devices;
devices.resize(ids.size());
auto lift = [](cl_device_id ptr) {
return detail::raw_device_ptr{ptr, false};
};
transform(ids.begin(), ids.end(), devices.begin(), lift);
detail::raw_context_ptr context;
context.reset(v2get(CAF_CLF(clCreateContext), nullptr,
static_cast<unsigned>(ids.size()),
ids.data(), pfn_notify, nullptr),
false);
vector<device_ptr> device_information;
for (auto& device_id : devices) {
device_information.push_back(device::create(context, device_id,
start_id++));
}
if (device_information.empty()) {
string errstr = "no devices for the platform found";
CAF_LOG_ERROR(CAF_ARG(errstr));
throw runtime_error(move(errstr));
}
auto name = platform_info(platform_id, CL_PLATFORM_NAME);
auto vendor = platform_info(platform_id, CL_PLATFORM_VENDOR);
auto version = platform_info(platform_id, CL_PLATFORM_VERSION);
return make_counted<platform>(platform_id, move(context), move(name),
move(vendor), move(version),
move(device_information));
}
string platform::platform_info(cl_platform_id platform_id,
unsigned info_flag) {
size_t size;
auto err = clGetPlatformInfo(platform_id, info_flag, 0, nullptr,
&size);
throwcl("clGetPlatformInfo", err);
vector<char> buffer(size);
v2callcl(CAF_CLF(clGetPlatformInfo), platform_id, info_flag,
sizeof(char) * size, buffer.data());
return string(buffer.data());
}
platform::platform(cl_platform_id platform_id, detail::raw_context_ptr context,
string name, string vendor, string version,
vector<device_ptr> devices)
: platform_id_(platform_id),
context_(std::move(context)),
name_(move(name)),
vendor_(move(vendor)),
version_(move(version)),
devices_(move(devices)) {
// nop
}
platform::~platform() {
// nop
}
} // namespace opencl
} // namespace caf
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| 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 <map>
#include <vector>
#include <string>
#include <cstring>
#include <iostream>
#include "caf/opencl/manager.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/opencl_err.hpp"
using namespace std;
namespace caf {
namespace opencl {
program::program(detail::raw_context_ptr context,
detail::raw_command_queue_ptr queue,
detail::raw_program_ptr prog,
map<string, detail::raw_kernel_ptr> available_kernels)
: context_(move(context)),
program_(move(prog)),
queue_(move(queue)),
available_kernels_(move(available_kernels)) {
// nop
}
program::~program() {
// nop
}
} // namespace opencl
} // namespace caf
#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/system_messages.hpp"
#include "caf/opencl/all.hpp"
using namespace std;
using namespace caf;
using namespace caf::opencl;
using caf::detail::tl_at;
using caf::detail::tl_head;
using caf::detail::type_list;
using caf::detail::limited_vector;
namespace {
using ivec = vector<int>;
using iref = mem_ref<int>;
using dims = opencl::dim_vec;
constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32;
constexpr size_t problem_size = 1024;
constexpr const char* kn_matrix = "matrix_square";
constexpr const char* kn_compiler_flag = "compiler_flag";
constexpr const char* kn_reduce = "reduce";
constexpr const char* kn_const = "const_mod";
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* kernel_source = R"__(
kernel void matrix_square(global const int* restrict matrix,
global int* restrict 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;
}
// http://developer.amd.com/resources/documentation-articles/
// articles-whitepapers/opencl-optimization-case-study-simple-reductions
kernel void reduce(global const int* restrict buffer,
global int* restrict result) {
local int scratch[512];
int local_index = get_local_id(0);
scratch[local_index] = buffer[get_global_id(0)];
barrier(CLK_LOCAL_MEM_FENCE);
for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
if (local_index < offset) {
int other = scratch[local_index + offset];
int mine = scratch[local_index];
scratch[local_index] = (mine < other) ? mine : other;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if (local_index == 0)
result[get_group_id(0)] = scratch[0];
}
kernel void const_mod(constant int* restrict input,
global int* restrict output) {
size_t idx = get_global_id(0);
output[idx] = input[0];
}
kernel void times_two(global int* restrict values) {
size_t idx = get_global_id(0);
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>
template<size_t Size>
class square_matrix {
public:
using value_type = ivec::value_type;
static constexpr size_t num_elements = Size * Size;
template <class Inspector>
friend typename Inspector::result_type inspect(Inspector& f,
square_matrix& x) {
return f(meta::type_name("square_matrix"), x.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];
}
using const_iterator = typename ivec::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_ = move(new_data);
}
private:
ivec data_;
};
template <class T>
vector<T> make_iota_vector(size_t num_elements) {
vector<T> result;
result.resize(num_elements);
iota(result.begin(), result.end(), T{0});
return result;
}
template <size_t Size>
square_matrix<Size> make_iota_matrix() {
square_matrix<Size> result;
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 string& description,
const vector<T>& expected,
const vector<T>& result) {
auto cond = (expected == result);
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 < result.size(); ++i) {
cout << " " << result[i];
}
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) {
auto& mngr = sys.opencl_manager();
auto opt = mngr.find_device_if([](const device_ptr){ return true; });
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 expected1{ 56, 62, 68, 74,
152, 174, 196, 218,
248, 286, 324, 362,
344, 398, 452, 506};
auto w1 = mngr.spawn(prog, kn_matrix,
opencl::nd_range{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 >> wrong_msg
);
opencl::nd_range range2{dims{matrix_size, matrix_size}};
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, range2,
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"
" (kernel passed directly)",
expected1, result);
}, others >> wrong_msg
);
const matrix_type expected2(move(expected1));
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)});
};
opencl::nd_range range3{dims{matrix_size, matrix_size}};
// let the runtime choose the device
auto w3 = mngr.spawn(mngr.create_program(kernel_source), kn_matrix, range3,
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 >> wrong_msg
);
opencl::nd_range range4{dims{matrix_size, matrix_size}};
auto w4 = mngr.spawn(prog, kn_matrix, range4,
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 >> wrong_msg
);
CAF_MESSAGE("Expecting exception (compiling invalid kernel, "
"semicolon is missing).");
try {
/* auto expected_error = */ mngr.create_program(kernel_source_error);
} catch (const exception& exc) {
std::string starts_with("clBuildProgram: CL_BUILD_PROGRAM_FAILURE");
auto cond = (strncmp(exc.what(), starts_with.c_str(),
starts_with.size()) == 0);
CAF_CHECK(cond);
if (!cond)
CAF_ERROR("Wrong exception cought for program build failure.");
}
// create program with opencl compiler flags
auto prog5 = mngr.create_program(kernel_source_compiler_flag, compiler_flag);
opencl::nd_range range5{dims{array_size}};
auto w5 = mngr.spawn(prog5, kn_compiler_flag, range5,
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 >> wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto max_wg_size = min(dev->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::nd_range range6{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, range6,
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
);
// calculator function for getting the size of the output
auto result_size_7 = [](const ivec&) {
return problem_size;
};
// constant memory arguments
const ivec arr7{problem_size};
auto w7 = mngr.spawn(kernel_source, kn_const,
opencl::nd_range{dims{problem_size}},
opencl::in<int>{},
opencl::out<int>{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 >> wrong_msg
);
}
void test_arguments(actor_system& sys) {
auto& mngr = sys.opencl_manager();
auto opt = mngr.find_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::nd_range{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
);
ivec input9 = make_iota_vector<int>(problem_size);
ivec expected9{input9};
for_each(begin(expected9), end(expected9), [](int& val){ val *= 2; });
auto w9 = mngr.spawn(kernel_source, kn_inout,
nd_range{dims{problem_size}},
opencl::in_out<int>{});
self->send(w9, move(input9));
self->receive(
[&](const ivec& result) {
check_vector_results("Testing in_out arugment", expected9, result);
}, others >> wrong_msg
);
ivec input10 = make_iota_vector<int>(problem_size);
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,
nd_range{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,
nd_range{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,
nd_range{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,
nd_range{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(opencl_basics) {
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();
}
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.find_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 = ::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 ::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 = ::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 ::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 = ::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 ::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 = ::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 ::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 ::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 ::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 ::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 ::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 ::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.find_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::nd_range{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);
nd_range range2{dims{array_size}};
auto w4 = mngr.spawn(prog2, kn_compiler_flag, range2,
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->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; });
nd_range range3{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, range3,
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,
nd_range{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.find_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 range = opencl::nd_range{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, range, 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, range,
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->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; });
nd_range range3{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, range3,
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.find_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 range = opencl::nd_range{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, range, 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, range,
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->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; });
nd_range range3{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, range3,
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.find_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 range = opencl::nd_range{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, range,
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, range,
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->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; });
nd_range range3{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, range3,
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.find_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;
nd_range range{dims{size}};
auto input1 = make_iota_vector<int>(size);
auto input2 = dev->global_argument(input1);
auto w1 = mngr.spawn(prog, kn_varying, range,
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, range,
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.find_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 range = nd_range{dims{problem_size}};
auto w1 = mngr.spawn(kernel_source, kn_inout, range,
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, range,
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, range,
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, range,
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.find_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
nd_range range{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, range,
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, range,
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.find_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 range = nd_range{dims{global_size}, {}, dims{local_size}};
auto w = mngr.spawn(kernel_source, kn_local, range,
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, range,
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(actor_facade) {
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