Commit aeef3fc0 authored by Dominik Charousset's avatar Dominik Charousset

Remove calls to deprecated APIs

parent 4356a5d4
...@@ -29,7 +29,6 @@ ...@@ -29,7 +29,6 @@
#include "caf/channel.hpp" #include "caf/channel.hpp"
#include "caf/to_string.hpp" #include "caf/to_string.hpp"
#include "cppa/tuple_cast.hpp"
#include "caf/intrusive_ptr.hpp" #include "caf/intrusive_ptr.hpp"
#include "caf/detail/int_list.hpp" #include "caf/detail/int_list.hpp"
...@@ -47,26 +46,28 @@ namespace opencl { ...@@ -47,26 +46,28 @@ namespace opencl {
class opencl_metainfo; class opencl_metainfo;
template <typename Signature> template <class Signature>
class actor_facade; class actor_facade;
template <typename Ret, typename... Args> template <class Ret, typename... Args>
class actor_facade<Ret(Args...)> : public abstract_actor { class actor_facade<Ret(Args...)> : public abstract_actor {
friend class command<actor_facade, Ret>; friend class command<actor_facade, Ret>;
public: public:
using args_tuple = using arg_types = detail::type_list<typename std::decay<Args>::type...>;
cow_tuple<typename std::decay<Args>::type...>; using arg_mapping = std::function<optional<message> (message&)>;
using arg_mapping = std::function<optional<args_tuple>(message)>;
using result_mapping = std::function<message(Ret&)>; using result_mapping = std::function<message(Ret&)>;
using evnt_vec = std::vector<cl_event>;
using args_vec = std::vector<mem_ptr>;
using command_type = command<actor_facade, Ret>;
static intrusive_ptr<actor_facade> static intrusive_ptr<actor_facade>
create(const program& prog, const char* kernel_name, arg_mapping map_args, create(const program& prog, const char* kernel_name,
result_mapping map_result, const dim_vec& global_dims, const dim_vec& global_dims, const dim_vec& offsets,
const dim_vec& offsets, const dim_vec& local_dims, const dim_vec& local_dims, size_t result_size,
size_t result_size) { arg_mapping map_args = arg_mapping{},
result_mapping map_result = result_mapping{}) {
if (global_dims.empty()) { if (global_dims.empty()) {
auto str = "OpenCL kernel needs at least 1 global dimension."; auto str = "OpenCL kernel needs at least 1 global dimension.";
CAF_LOGM_ERROR("caf::opencl::actor_facade", str); CAF_LOGM_ERROR("caf::opencl::actor_facade", str);
...@@ -84,32 +85,50 @@ class actor_facade<Ret(Args...)> : public abstract_actor { ...@@ -84,32 +85,50 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
check_vec(offsets, "offsets"); check_vec(offsets, "offsets");
check_vec(local_dims, "local dimensions"); check_vec(local_dims, "local dimensions");
kernel_ptr kernel; kernel_ptr kernel;
kernel.adopt(v2get(CAF_CLF(clCreateKernel), prog.m_program.get(), kernel.reset(v2get(CAF_CLF(clCreateKernel), prog.m_program.get(),
kernel_name)); kernel_name),
false);
if (result_size == 0) { if (result_size == 0) {
result_size = std::accumulate(global_dims.begin(), global_dims.end(), result_size = std::accumulate(global_dims.begin(), global_dims.end(),
size_t{1}, std::multiplies<size_t>{}); size_t{1}, std::multiplies<size_t>{});
} }
return new actor_facade<Ret(Args...)>{ return new actor_facade(prog, kernel, global_dims, offsets, local_dims,
prog, kernel, global_dims, offsets, result_size, std::move(map_args),
local_dims, std::move(map_args), std::move(map_result), result_size}; std::move(map_result));
} }
void enqueue(const actor_addr &sender, message_id mid, message content, void enqueue(const actor_addr &sender, message_id mid, message content,
execution_unit*) override { execution_unit*) override {
CAF_LOG_TRACE(""); CAF_LOG_TRACE("");
typename detail::il_indices<detail::type_list<Args...>>::type indices; if (m_map_args) {
enqueue_impl(sender, mid, std::move(content), indices); auto mapped = m_map_args(content);
if (!mapped) {
return;
}
content = std::move(*mapped);
}
typename detail::il_indices<arg_types>::type indices;
if (!content.match_elements(arg_types{})) {
return;
}
response_promise handle{this->address(), sender, mid.response_id()};
evnt_vec events;
args_vec arguments;
add_arguments_to_kernel<Ret>(events, arguments, m_result_size,
content, indices);
auto cmd = detail::make_counted<command_type>(handle, this,
std::move(events),
std::move(arguments),
m_result_size,
std::move(content));
cmd->enqueue();
} }
private: private:
using evnt_vec = std::vector<cl_event>;
using args_vec = std::vector<mem_ptr>;
actor_facade(const program& prog, kernel_ptr kernel, actor_facade(const program& prog, kernel_ptr kernel,
const dim_vec& global_dimensions, const dim_vec& global_offsets, const dim_vec& global_dimensions, const dim_vec& global_offsets,
const dim_vec& local_dimensions, arg_mapping map_args, const dim_vec& local_dimensions, size_t result_size,
result_mapping map_result, size_t result_size) arg_mapping map_args, result_mapping map_result)
: m_kernel(kernel), : m_kernel(kernel),
m_program(prog.m_program), m_program(prog.m_program),
m_context(prog.m_context), m_context(prog.m_context),
...@@ -117,45 +136,16 @@ class actor_facade<Ret(Args...)> : public abstract_actor { ...@@ -117,45 +136,16 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
m_global_dimensions(global_dimensions), m_global_dimensions(global_dimensions),
m_global_offsets(global_offsets), m_global_offsets(global_offsets),
m_local_dimensions(local_dimensions), m_local_dimensions(local_dimensions),
m_result_size(result_size),
m_map_args(std::move(map_args)), m_map_args(std::move(map_args)),
m_map_result(std::move(map_result)), m_map_result(std::move(map_result)) {
m_result_size(result_size) {
CAF_LOG_TRACE("id: " << this->id()); CAF_LOG_TRACE("id: " << this->id());
} }
template <long... Is> void add_arguments_to_kernel_rec(evnt_vec&, args_vec& arguments, message&,
void enqueue_impl(const actor_addr& sender, message_id mid, message msg, detail::int_list<>) {
detail::int_list<Is...>) {
auto opt = m_map_args(std::move(msg));
if (opt) {
response_promise handle{this->address(), sender, mid.response_id()};
evnt_vec events;
args_vec arguments;
add_arguments_to_kernel<Ret>(events, arguments, m_result_size,
get_ref<Is>(*opt)...);
auto cmd = detail::make_counted<command<actor_facade, Ret>>(
handle, this, std::move(events), std::move(arguments), m_result_size,
*opt);
cmd->enqueue();
} else {
CAF_LOGMF(CAF_ERROR, "actor_facade::enqueue() tuple_cast failed.");
}
}
kernel_ptr m_kernel;
program_ptr m_program;
context_ptr m_context;
command_queue_ptr m_queue;
dim_vec m_global_dimensions;
dim_vec m_global_offsets;
dim_vec m_local_dimensions;
arg_mapping m_map_args;
result_mapping m_map_result;
size_t m_result_size;
void add_arguments_to_kernel_rec(evnt_vec&, args_vec& arguments) {
// rotate left (output buffer to the end) // rotate left (output buffer to the end)
rotate(begin(arguments), begin(arguments) + 1, end(arguments)); std::rotate(arguments.begin(), arguments.begin() + 1, arguments.end());
for (cl_uint i = 0; i < arguments.size(); ++i) { for (cl_uint i = 0; i < arguments.size(); ++i) {
v1callcl(CAF_CLF(clSetKernelArg), m_kernel.get(), i, v1callcl(CAF_CLF(clSetKernelArg), m_kernel.get(), i,
sizeof(cl_mem), static_cast<void*>(&arguments[i])); sizeof(cl_mem), static_cast<void*>(&arguments[i]));
...@@ -163,34 +153,48 @@ class actor_facade<Ret(Args...)> : public abstract_actor { ...@@ -163,34 +153,48 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
clFlush(m_queue.get()); clFlush(m_queue.get());
} }
template <typename T0, typename... Ts> template <long I, long... Is>
void add_arguments_to_kernel_rec(evnt_vec& events, args_vec& arguments, void add_arguments_to_kernel_rec(evnt_vec& events, args_vec& arguments,
T0& arg0, Ts&... args) { message& msg, detail::int_list<I, Is...>) {
size_t buffer_size = sizeof(typename T0::value_type) * arg0.size(); using value_type = typename detail::tl_at<arg_types, I>::type;
auto& arg = msg.get_as<value_type>(I);
size_t buffer_size = sizeof(value_type) * arg.size();
auto buffer = v2get(CAF_CLF(clCreateBuffer), m_context.get(), auto buffer = v2get(CAF_CLF(clCreateBuffer), m_context.get(),
cl_mem_flags{CL_MEM_READ_ONLY}, buffer_size, nullptr); cl_mem_flags{CL_MEM_READ_ONLY}, buffer_size, nullptr);
cl_event event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer), cl_event event = v1get<cl_event>(CAF_CLF(clEnqueueWriteBuffer),
m_queue.get(), buffer, cl_bool{CL_FALSE}, m_queue.get(), buffer, cl_bool{CL_FALSE},
cl_uint{0},buffer_size, arg0.data()); cl_uint{0}, buffer_size, arg.data());
events.push_back(std::move(event)); events.push_back(std::move(event));
mem_ptr tmp; mem_ptr tmp;
tmp.adopt(std::move(buffer)); tmp.reset(buffer, false);
arguments.push_back(tmp); arguments.push_back(tmp);
add_arguments_to_kernel_rec(events, arguments, args...); add_arguments_to_kernel_rec(events, arguments, msg,
detail::int_list<Is...>{});
} }
template <typename R, typename... Ts> template <class R, class Token>
void add_arguments_to_kernel(evnt_vec& events, args_vec& arguments, void add_arguments_to_kernel(evnt_vec& events, args_vec& arguments,
size_t ret_size, Ts&&... args) { size_t ret_size, message& msg, Token tk) {
arguments.clear(); arguments.clear();
auto buf = v2get(CAF_CLF(clCreateBuffer), m_context.get(), auto buf = v2get(CAF_CLF(clCreateBuffer), m_context.get(),
cl_mem_flags{CL_MEM_WRITE_ONLY}, cl_mem_flags{CL_MEM_WRITE_ONLY},
sizeof(typename R::value_type) * ret_size, nullptr); sizeof(typename R::value_type) * ret_size, nullptr);
mem_ptr tmp; mem_ptr tmp;
tmp.adopt(std::move(buf)); tmp.reset(buf, false);
arguments.push_back(tmp); arguments.push_back(tmp);
add_arguments_to_kernel_rec(events, arguments, std::forward<Ts>(args)...); add_arguments_to_kernel_rec(events, arguments, msg, tk);
} }
kernel_ptr m_kernel;
program_ptr m_program;
context_ptr m_context;
command_queue_ptr m_queue;
dim_vec m_global_dimensions;
dim_vec m_global_offsets;
dim_vec m_local_dimensions;
size_t m_result_size;
arg_mapping m_map_args;
result_mapping m_map_result;
}; };
} // namespace opencl } // namespace opencl
......
...@@ -25,7 +25,6 @@ ...@@ -25,7 +25,6 @@
#include "caf/actor_cast.hpp" #include "caf/actor_cast.hpp"
#include "caf/optional.hpp" #include "caf/optional.hpp"
#include "cppa/cow_tuple.hpp"
#include "caf/detail/limited_vector.hpp" #include "caf/detail/limited_vector.hpp"
...@@ -48,127 +47,98 @@ struct carr_to_vec<T*> { ...@@ -48,127 +47,98 @@ struct carr_to_vec<T*> {
using type = std::vector<T>; using type = std::vector<T>;
}; };
template <typename Signature, typename SecondSignature = void> template <class Signature>
struct cl_spawn_helper; struct cl_spawn_helper;
template <typename R, typename... Ts> template <typename R, typename... Ts>
struct cl_spawn_helper<R(Ts...), void> { struct cl_spawn_helper<R(Ts...)> {
using res_t = typename carr_to_vec<R>::type;
using result_type = typename carr_to_vec<R>::type; using impl = opencl::actor_facade<res_t (typename carr_to_vec<Ts>::type...)>;
using map_arg_fun = std::function<optional<message> (message&)>;
using impl = opencl::actor_facade<
result_type(typename carr_to_vec<typename carr_to_vec<Ts>::type>::type...)>;
using map_arg_fun = typename impl::arg_mapping;
using map_res_fun = typename impl::result_mapping; using map_res_fun = typename impl::result_mapping;
template <typename... Us> template <typename... Us>
actor operator()(map_arg_fun f0, map_res_fun f1, const opencl::program& p, actor operator()(const opencl::program& p, const char* fn, Us&&... vs) const {
const char* fname, Us&&... args) const { return actor_cast<actor>(impl::create(p, fn, std::forward<Us>(vs)...));
using std::move;
using std::forward;
return actor_cast<actor>(
impl::create(p, fname, move(f0), move(f1), forward<Us>(args)...));
}
template <typename... Us>
actor operator()(const opencl::program& p, const char* fname,
Us&&... args) const {
using std::move;
using std::forward;
map_arg_fun f0 = [](message msg) {
return tuple_cast<typename std::decay<
typename carr_to_vec<Ts>::type>::type...>(msg);
};
map_res_fun f1 = [](result_type& result) {
return make_message(move(result));
};
return actor_cast<actor>(
impl::create(p, fname, move(f0), move(f1), forward<Us>(args)...));
} }
}; };
template <typename R, typename... Ts>
struct cl_spawn_helper<std::function<optional<cow_tuple<Ts...>>(message)>,
std::function<message(R&)>>
: cl_spawn_helper<R(Ts...)> {};
} // namespace detail } // namespace detail
/** /**
* @brief Creates a new actor facade for an OpenCL kernel that invokes * Creates a new actor facade for an OpenCL kernel that invokes
* the function named @p fname from @p prog. * the function named `fname` from `prog`.
* @throws std::runtime_error if more than three dimensions are set, * @throws std::runtime_error if more than three dimensions are set,
* <tt>dims.empty()</tt>, or @p clCreateKernel * `dims.empty()`, or `clCreateKernel` failed.
* failed.
*/ */
template <typename Signature, typename... Ts> template <class Signature>
inline actor actor spawn_cl(const opencl::program& prog,
spawn_cl(const opencl::program& prog, const char* fname, const char* fname,
const opencl::dim_vec& dims, const opencl::dim_vec& offset = {}, const opencl::dim_vec& dims,
const opencl::dim_vec& local_dims = {}, size_t result_size = 0) { const opencl::dim_vec& offset = {},
using std::move; const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
detail::cl_spawn_helper<Signature> f; detail::cl_spawn_helper<Signature> f;
return f(prog, fname, dims, offset, local_dims, result_size); return f(prog, fname, dims, offset, local_dims, result_size);
} }
/** /**
* @brief Compiles @p source and creates a new actor facade for an OpenCL kernel * Compiles `source` and creates a new actor facade for an OpenCL kernel
* that invokes the function named @p fname. * that invokes the function named `fname`.
* @throws std::runtime_error if more than three dimensions are set, * @throws std::runtime_error if more than three dimensions are set,
* <tt>dims.empty()</tt>, a compilation error * <tt>dims.empty()</tt>, a compilation error
* occured, or @p clCreateKernel failed. * occured, or @p clCreateKernel failed.
*/ */
template <typename Signature, typename... Ts> template <class Signature>
inline actor actor spawn_cl(const char* source,
spawn_cl(const char* source, const char* fname, const opencl::dim_vec& dims, const char* fname,
const opencl::dim_vec& offset = {}, const opencl::dim_vec& dims,
const opencl::dim_vec& local_dims = {}, size_t result_size = 0) { const opencl::dim_vec& offset = {},
using std::move; const opencl::dim_vec& local_dims = {},
return spawn_cl<Signature, Ts...>(opencl::program::create(source), fname, size_t result_size = 0) {
dims, offset, local_dims, result_size); return spawn_cl<Signature>(opencl::program::create(source), fname,
dims, offset, local_dims, result_size);
} }
/** /**
* @brief Creates a new actor facade for an OpenCL kernel that invokes * Creates a new actor facade for an OpenCL kernel that invokes
* the function named @p fname from @p prog, using @p map_args * the function named `fname` from `prog`.
* to extract the function arguments from incoming messages and
* @p map_result to transform the result before sending it as response.
* @throws std::runtime_error if more than three dimensions are set, * @throws std::runtime_error if more than three dimensions are set,
* <tt>dims.empty()</tt>, or @p clCreateKernel * `dims.empty()`, or `clCreateKernel` failed.
* failed.
*/ */
template <typename MapArgs, typename MapResult> template <class Signature, class Fun>
inline actor actor spawn_cl(const opencl::program& prog,
spawn_cl(const opencl::program& prog, const char* fname, MapArgs map_args, const char* fname,
MapResult map_result, const opencl::dim_vec& dims, std::function<optional<message> (message&)> map_args,
const opencl::dim_vec& offset = {}, Fun map_result,
const opencl::dim_vec& local_dims = {}, size_t result_size = 0) { const opencl::dim_vec& dims,
using std::move; const opencl::dim_vec& offset = {},
using f0 = typename detail::get_callable_trait<MapArgs>::fun_type; const opencl::dim_vec& local_dims = {},
using f1 = typename detail::get_callable_trait<MapResult>::fun_type; size_t result_size = 0) {
detail::cl_spawn_helper<f0, f1> f; detail::cl_spawn_helper<Signature> f;
return f(f0{move(map_args)}, f1{move(map_result)}, prog, fname, dims, offset, return f(prog, fname, dims, offset, local_dims, result_size,
local_dims, result_size); std::move(map_args), std::move(map_result));
} }
/** /**
* @brief Compiles @p source and creates a new actor facade for an OpenCL kernel * Compiles `source` and creates a new actor facade for an OpenCL kernel
* that invokes the function named @p fname, using @p map_args * that invokes the function named `fname`.
* to extract the function arguments from incoming messages and
* @p map_result to transform the result before sending it as response.
* @throws std::runtime_error if more than three dimensions are set, * @throws std::runtime_error if more than three dimensions are set,
* <tt>dims.empty()</tt>, a compilation error * <tt>dims.empty()</tt>, a compilation error
* occured, or @p clCreateKernel failed. * occured, or @p clCreateKernel failed.
*/ */
template <typename MapArgs, typename MapResult> template <class Signature, class Fun>
inline actor actor spawn_cl(const char* source,
spawn_cl(const char* source, const char* fun_name, MapArgs map_args, const char* fname,
MapResult map_result, const opencl::dim_vec& dims, std::function<optional<message> (message&)> map_args,
const opencl::dim_vec& offset = {}, Fun map_result,
const opencl::dim_vec& local_dims = {}, size_t result_size = 0) { const opencl::dim_vec& dims,
using std::move; const opencl::dim_vec& offset = {},
return spawn_cl(opencl::program::create(source), fun_name, move(map_args), const opencl::dim_vec& local_dims = {},
move(map_result), dims, offset, local_dims, result_size); size_t result_size = 0) {
detail::cl_spawn_helper<Signature> f;
return f(opencl::program::create(source), fname, dims, offset, local_dims,
result_size, std::move(map_args), std::move(map_result));
} }
} // namespace caf } // namespace caf
......
...@@ -142,28 +142,31 @@ void multiplier(event_based_actor* self) { ...@@ -142,28 +142,31 @@ void multiplier(event_based_actor* self) {
cout << "calculating square of matrix:" << endl cout << "calculating square of matrix:" << endl
<< to_string(m1) << endl; << to_string(m1) << endl;
auto unbox_args = [](message& msg) -> optional<message> {
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 // spawn an opencl actor
// 1st arg: source code of one or more opencl kernels // 1st arg: source code of one or more opencl kernels
// 2nd arg: name of the kernel to use // 2nd arg: name of the kernel to use
auto worker = spawn_cl(kernel_source, kernel_name, auto worker = spawn_cl<fvec (fvec&, fvec&)>(kernel_source, kernel_name,
// 3rd arg: the opencl function operates on vectors, // 3rd arg: the opencl function operates on vectors,
// this function converts a tuple of two matrices // this function converts a tuple of two matrices
// to a tuple of vectors; note that this function returns // to a tuple of vectors; note that this function returns
// an option (an empty results causes the actor to ignore // an option (an empty results causes the actor to ignore
// the message) // the message)
[] (message msg) -> optional<cow_tuple<fvec, fvec>> { unbox_args,
auto opt = tuple_cast<matrix_type, matrix_type>(msg);
if (opt) {
return make_cow_tuple(move(get_ref<0>(*opt).data()),
move(get_ref<1>(*opt).data()));
}
return none;
},
// 4th arg: converts the ouptut vector back to a matrix that is then // 4th arg: converts the ouptut vector back to a matrix that is then
// used as response message // used as response message
[] (fvec& result) -> message { box_res,
return make_message(matrix_type{move(result)});
},
// 5th arg: global dimension arguments for opencl's enqueue, // 5th arg: global dimension arguments for opencl's enqueue,
// creates matrix_size * matrix_size global work items // creates matrix_size * matrix_size global work items
{matrix_size, matrix_size} {matrix_size, matrix_size}
......
...@@ -9,14 +9,13 @@ ...@@ -9,14 +9,13 @@
#include "caf/all.hpp" #include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp" #include "caf/opencl/spawn_cl.hpp"
using namespace std;
using namespace caf; using namespace caf;
using namespace caf::opencl; using namespace caf::opencl;
namespace { namespace {
using ivec = vector<int>; using ivec = std::vector<int>;
using fvec = vector<float>; using fvec = std::vector<float>;
constexpr size_t matrix_size = 4; constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32; constexpr size_t array_size = 32;
...@@ -28,11 +27,11 @@ constexpr const char* kernel_name_compiler_flag = "compiler_flag"; ...@@ -28,11 +27,11 @@ constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_reduce = "reduce"; constexpr const char* kernel_name_reduce = "reduce";
constexpr const char* kernel_name_const = "const_mod"; constexpr const char* kernel_name_const = "const_mod";
constexpr const char* compiler_flag = "-D OPENCL_CPPA_TEST_FLAG"; constexpr const char* compiler_flag = "-D CAF_OPENCL_TEST_FLAG";
constexpr const char* kernel_source = R"__( constexpr const char* kernel_source = R"__(
__kernel void matrix_square(__global int* matrix, __kernel void matrix_square(__global int* matrix,
__global int* output) { __global int* output) {
size_t size = get_global_size(0); // == get_global_size_(1); size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0); size_t x = get_global_id(0);
size_t y = get_global_id(1); size_t y = get_global_id(1);
...@@ -52,13 +51,13 @@ constexpr const char* kernel_source_error = R"__( ...@@ -52,13 +51,13 @@ constexpr const char* kernel_source_error = R"__(
constexpr const char* kernel_source_compiler_flag = R"__( constexpr const char* kernel_source_compiler_flag = R"__(
__kernel void compiler_flag(__global int* input, __kernel void compiler_flag(__global int* input,
__global int* output) { __global int* output) {
size_t x = get_global_id(0); size_t x = get_global_id(0);
#ifdef OPENCL_CPPA_TEST_FLAG # ifdef CAF_OPENCL_TEST_FLAG
output[x] = input[x]; output[x] = input[x];
#else # else
output[x] = 0; output[x] = 0;
#endif # endif
} }
)__"; )__";
...@@ -66,12 +65,11 @@ constexpr const char* kernel_source_compiler_flag = R"__( ...@@ -66,12 +65,11 @@ constexpr const char* kernel_source_compiler_flag = R"__(
// opencl-optimization-case-study-simple-reductions // opencl-optimization-case-study-simple-reductions
constexpr const char* kernel_source_reduce = R"__( constexpr const char* kernel_source_reduce = R"__(
__kernel void reduce(__global int* buffer, __kernel void reduce(__global int* buffer,
__global int* result) { __global int* result) {
__local int scratch[512]; __local int scratch[512];
int local_index = get_local_id(0); int local_index = get_local_id(0);
scratch[local_index] = buffer[get_global_id(0)]; scratch[local_index] = buffer[get_global_id(0)];
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) { for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
if (local_index < offset) { if (local_index < offset) {
int other = scratch[local_index + offset]; int other = scratch[local_index + offset];
...@@ -88,19 +86,17 @@ constexpr const char* kernel_source_reduce = R"__( ...@@ -88,19 +86,17 @@ constexpr const char* kernel_source_reduce = R"__(
constexpr const char* kernel_source_const = R"__( constexpr const char* kernel_source_const = R"__(
__kernel void const_mod(__constant int* input, __kernel void const_mod(__constant int* input,
__global int* output) { __global int* output) {
size_t idx = get_global_id(0); size_t idx = get_global_id(0);
output[idx] = input[0]; output[idx] = input[0];
} }
)__"; )__";
} } // namespace <anonymous>
template<size_t Size> template<size_t Size>
class square_matrix { class square_matrix {
public: public:
static constexpr size_t num_elements = Size * Size; static constexpr size_t num_elements = Size * Size;
static void announce() { static void announce() {
...@@ -112,51 +108,74 @@ class square_matrix { ...@@ -112,51 +108,74 @@ class square_matrix {
square_matrix& operator=(square_matrix&&) = default; square_matrix& operator=(square_matrix&&) = default;
square_matrix& operator=(const square_matrix&) = default; square_matrix& operator=(const square_matrix&) = default;
square_matrix() : m_data(num_elements) { } square_matrix() : m_data(num_elements) {
// nop
}
explicit square_matrix(ivec d) : m_data(move(d)) { explicit square_matrix(ivec d) : m_data(move(d)) {
assert(m_data.size() == num_elements); assert(m_data.size() == num_elements);
} }
inline float& operator()(size_t column, size_t row) { float& operator()(size_t column, size_t row) {
return m_data[column + row * Size]; return m_data[column + row * Size];
} }
inline const float& operator()(size_t column, size_t row) const { const float& operator()(size_t column, size_t row) const {
return m_data[column + row * Size]; return m_data[column + row * Size];
} }
inline void iota_fill() {
iota(m_data.begin(), m_data.end(), 0);
}
typedef typename ivec::const_iterator const_iterator; typedef typename ivec::const_iterator const_iterator;
const_iterator begin() const { return m_data.begin(); } const_iterator begin() const {
return m_data.begin();
}
const_iterator end() const { return m_data.end(); } const_iterator end() const {
return m_data.end();
}
ivec& data() { return m_data; } ivec& data() {
return m_data;
}
const ivec& data() const { return m_data; } const ivec& data() const {
return m_data;
}
void data(ivec new_data) { m_data = std::move(new_data); } void data(ivec new_data) {
m_data = std::move(new_data);
}
private: private:
ivec m_data; ivec m_data;
}; };
template <class T>
std::vector<T> make_iota_vector(size_t num_elements) {
std::vector<T> result;
result.resize(num_elements);
std::iota(result.begin(), result.end(), T{0});
return result;
}
template <size_t Size>
square_matrix<Size> make_iota_matrix() {
square_matrix<Size> result;
std::iota(result.data().begin(), result.data().end(), 0);
return result;
}
template<size_t Size> template<size_t Size>
inline bool operator==(const square_matrix<Size>& lhs, bool operator==(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) { const square_matrix<Size>& rhs) {
return equal(lhs.begin(), lhs.end(), rhs.begin()); return lhs.data() == rhs.data();
} }
template<size_t Size> template<size_t Size>
inline bool operator!=(const square_matrix<Size>& lhs, bool operator!=(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) { const square_matrix<Size>& rhs) {
return !(lhs == rhs); return !(lhs == rhs);
} }
...@@ -170,154 +189,117 @@ size_t get_max_workgroup_size(size_t device_id, size_t dimension) { ...@@ -170,154 +189,117 @@ size_t get_max_workgroup_size(size_t device_id, size_t dimension) {
} }
void test_opencl() { void test_opencl() {
scoped_actor self; scoped_actor self;
const ivec expected1{ 56, 62, 68, 74,
const ivec expected1{ 56, 62, 68, 74 152, 174, 196, 218,
, 152, 174, 196, 218 248, 286, 324, 362,
, 248, 286, 324, 362 344, 398, 452, 506};
, 344, 398, 452, 506}; auto w1 = spawn_cl<ivec (ivec&)>(program::create(kernel_source),
kernel_name,
auto worker1 = spawn_cl<ivec(ivec&)>(program::create(kernel_source), {matrix_size, matrix_size});
kernel_name, self->send(w1, make_iota_vector<int>(matrix_size * matrix_size));
{matrix_size, matrix_size});
ivec m1(matrix_size * matrix_size);
iota(begin(m1), end(m1), 0);
self->send(worker1, move(m1));
self->receive ( self->receive (
on_arg_match >> [&] (const ivec& result) { [&](const ivec& result) {
CAF_CHECK(equal(begin(expected1), end(expected1), begin(result))); CAF_CHECK(result == expected1);
} }
); );
auto w2 = spawn_cl<ivec (ivec&)>(kernel_source, kernel_name,
auto worker2 = spawn_cl<ivec(ivec&)>(kernel_source, {matrix_size, matrix_size});
kernel_name, self->send(w2, make_iota_vector<int>(matrix_size * matrix_size));
{matrix_size, matrix_size});
ivec m2(matrix_size * matrix_size);
iota(begin(m2), end(m2), 0);
self->send(worker2, move(m2));
self->receive ( self->receive (
on_arg_match >> [&] (const ivec& result) { [&](const ivec& result) {
CAF_CHECK(equal(begin(expected1), end(expected1), begin(result))); CAF_CHECK(result == expected1);
} }
); );
const matrix_type expected2(std::move(expected1));
const matrix_type expected2(move(expected1)); auto map_arg = [](message& msg) -> optional<message> {
return msg.apply(
auto map_args = [] (message msg) -> optional<cow_tuple<ivec>> { [](matrix_type& mx) {
auto opt = tuple_cast<matrix_type>(msg); return make_message(std::move(mx.data()));
if (opt) { }
return make_cow_tuple(move(get_ref<0>(*opt).data())); );
}
return none;
}; };
auto map_res = [](ivec& result) -> message {
auto map_results = [] (ivec& result) -> message { return make_message(matrix_type{std::move(result)});
return make_message(matrix_type{move(result)});
}; };
auto w3 = spawn_cl<ivec (ivec&)>(program::create(kernel_source), kernel_name,
matrix_type m3; map_arg, map_res, {matrix_size, matrix_size});
m3.iota_fill(); self->send(w3, make_iota_matrix<matrix_size>());
auto worker3 = spawn_cl(program::create(kernel_source), kernel_name,
map_args, map_results,
{matrix_size, matrix_size});
self->send(worker3, move(m3));
self->receive ( self->receive (
on_arg_match >> [&] (const matrix_type& result) { [&](const matrix_type& result) {
CAF_CHECK(expected2 == result); CAF_CHECK(expected2 == result);
} }
); );
auto w4 = spawn_cl<ivec (ivec&)>(kernel_source, kernel_name,
matrix_type m4; map_arg, map_res,
m4.iota_fill(); {matrix_size, matrix_size});
auto worker4 = spawn_cl(kernel_source, kernel_name, self->send(w4, make_iota_matrix<matrix_size>());
map_args, map_results,
{matrix_size, matrix_size}
);
self->send(worker4, move(m4));
self->receive ( self->receive (
on_arg_match >> [&] (const matrix_type& result) { [&](const matrix_type& result) {
CAF_CHECK(expected2 == result); CAF_CHECK(expected2 == result);
} }
); );
try { try {
program create_error = program::create(kernel_source_error); auto create_error = program::create(kernel_source_error);
} }
catch (const exception& exc) { catch (const std::exception& exc) {
cout << exc.what() << endl; CAF_PRINT(exc.what());
CAF_CHECK_EQUAL("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what()); CAF_CHECK_EQUAL("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what());
} }
// test for opencl compiler flags // test for opencl compiler flags
ivec arr5(array_size);
iota(begin(arr5), end(arr5), 0);
auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag); auto prog5 = program::create(kernel_source_compiler_flag, compiler_flag);
auto worker5 = spawn_cl<ivec(ivec&)>(prog5, kernel_name_compiler_flag, {array_size}); auto w5 = spawn_cl<ivec (ivec&)>(prog5, kernel_name_compiler_flag, {array_size});
self->send(worker5, move(arr5)); self->send(w5, make_iota_vector<int>(array_size));
auto expected3 = make_iota_vector<int>(array_size);
ivec expected3(array_size);
iota(begin(expected3), end(expected3), 0);
self->receive ( self->receive (
on_arg_match >> [&] (const ivec& result) { [&](const ivec& result) {
CAF_CHECK(equal(begin(expected3), end(expected3), begin(result))); CAF_CHECK(result == expected3);
} }
); );
// test for manuel return size selection // test for manuel return size selection (max workgroup size 1d)
const int max_workgroup_size = static_cast<int>(get_max_workgroup_size(0,1)); // max workgroup size (1d) const int max_workgroup_size = static_cast<int>(get_max_workgroup_size(0,1));
const size_t reduce_buffer_size = static_cast<size_t>(max_workgroup_size * 8); const int reduce_buffer_size = max_workgroup_size * 8;
const size_t reduce_local_size = static_cast<size_t>(max_workgroup_size); const int reduce_local_size = max_workgroup_size;
const size_t reduce_work_groups = reduce_buffer_size / reduce_local_size; const int reduce_work_groups = reduce_buffer_size / reduce_local_size;
const size_t reduce_global_size = reduce_buffer_size; const int reduce_global_size = reduce_buffer_size;
const size_t reduce_result_size = reduce_work_groups; const int reduce_result_size = reduce_work_groups;
ivec arr6(static_cast<size_t>(reduce_buffer_size));
ivec arr6(reduce_buffer_size); int n = static_cast<int>(arr6.capacity());
int n{static_cast<int>(arr6.capacity())}; std::generate(arr6.begin(), arr6.end(), [&]{ return --n; });
generate(begin(arr6), end(arr6), [&]{ return --n; }); auto w6 = spawn_cl<ivec (ivec&)>(kernel_source_reduce, kernel_name_reduce,
auto worker6 = spawn_cl<ivec(ivec&)>(kernel_source_reduce, {static_cast<size_t>(reduce_global_size)},
kernel_name_reduce, {},
{reduce_global_size}, {static_cast<size_t>(reduce_local_size)},
{}, static_cast<size_t>(reduce_result_size));
{reduce_local_size}, self->send(w6, move(arr6));
reduce_result_size); ivec expected4{max_workgroup_size * 7, max_workgroup_size * 6,
self->send(worker6, move(arr6)); max_workgroup_size * 5, max_workgroup_size * 4,
const ivec expected4{ max_workgroup_size * 7, max_workgroup_size * 6, max_workgroup_size * 3, max_workgroup_size * 2,
max_workgroup_size * 5, max_workgroup_size * 4, max_workgroup_size, 0};
max_workgroup_size * 3, max_workgroup_size * 2, self->receive(
max_workgroup_size, 0 }; [&](const ivec& result) {
self->receive ( CAF_CHECK(result == expected4);
on_arg_match >> [&] (const ivec& result) {
CAF_CHECK(equal(begin(expected4), end(expected4), begin(result)));
} }
); );
// constant memory arguments // constant memory arguments
const ivec arr7{magic_number}; const ivec arr7{magic_number};
auto worker7 = spawn_cl<ivec(ivec&)>(kernel_source_const, auto w7 = spawn_cl<ivec (ivec&)>(kernel_source_const, kernel_name_const,
kernel_name_const, {magic_number});
{magic_number}); self->send(w7, move(arr7));
self->send(worker7, move(arr7));
ivec expected5(magic_number); ivec expected5(magic_number);
fill(begin(expected5), end(expected5), magic_number); fill(begin(expected5), end(expected5), magic_number);
self->receive( self->receive(
on_arg_match >> [&] (const ivec& result) { [&](const ivec& result) {
CAF_CHECK(equal(begin(expected5), end(expected5), begin(result))); CAF_CHECK(result == expected5);
} }
); );
} }
int main() { int main() {
CAF_TEST(tkest_opencl); CAF_TEST(test_opencl);
announce<ivec>("ivec");
matrix_type::announce();
test_opencl(); test_opencl();
await_all_actors_done(); await_all_actors_done();
shutdown(); shutdown();
return CAF_TEST_RESULT(); return CAF_TEST_RESULT();
} }
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