Commit f314c210 authored by Joseph Noir's avatar Joseph Noir

Rename spawn_config to nd_range

parent 522e26e3
......@@ -54,7 +54,7 @@ struct dummy_size_calculator {
}
};
/// Mark an a spawn_cl template argument as input only
/// Mark an a spawn template argument as input only
template <class Arg, class Tag = val>
struct in {
static_assert(std::is_same<Tag,val>::value || std::is_same<Tag,mref>::value,
......@@ -63,7 +63,7 @@ struct in {
using arg_type = typename std::decay<Arg>::type;
};
/// Mark an a spawn_cl template argument as input and output
/// Mark an a spawn template argument as input and output
template <class Arg, class TagIn = val, class TagOut = val>
struct in_out {
static_assert(
......
......@@ -37,7 +37,7 @@
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/spawn_config.hpp"
#include "caf/opencl/nd_range.hpp"
namespace caf {
namespace opencl {
......@@ -60,7 +60,7 @@ public:
std::vector<size_t> lengths,
message msg,
std::tuple<Ts...> output_tuple,
spawn_config config)
nd_range range)
: lengths_(std::move(lengths)),
promise_(std::move(promise)),
cl_actor_(std::move(parent)),
......@@ -70,7 +70,7 @@ public:
scratch_buffers_(std::move(scratches)),
results_(output_tuple),
msg_(std::move(msg)),
config_(std::move(config)) {
range_(std::move(range)) {
// nop
}
......@@ -107,10 +107,10 @@ public:
mem_out_events_.emplace_back();
cl_int err = clEnqueueNDRangeKernel(
parent->queue_.get(), parent->kernel_.get(),
static_cast<unsigned int>(config_.dimensions().size()),
data_or_nullptr(config_.offsets()),
data_or_nullptr(config_.dimensions()),
data_or_nullptr(config_.local_dimensions()),
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()
......@@ -183,10 +183,10 @@ public:
cl_event execution_event;
cl_int err = clEnqueueNDRangeKernel(
parent->queue_.get(), parent->kernel_.get(),
static_cast<cl_uint>(config_.dimensions().size()),
data_or_nullptr(config_.offsets()),
data_or_nullptr(config_.dimensions()),
data_or_nullptr(config_.local_dimensions()),
static_cast<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
......@@ -274,7 +274,7 @@ private:
std::vector<cl_mem_ptr> scratch_buffers_;
std::tuple<Ts...> results_;
message msg_; // keeps the argument buffers alive for async copy to device
spawn_config config_;
nd_range range_;
};
} // namespace opencl
......
......@@ -35,28 +35,28 @@ struct cl_spawn_helper {
using map_out_fun = typename impl::output_mapping;
actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::spawn_config& spawn_cfg,
const char* fn, const opencl::nd_range& range,
Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg),
p, fn, spawn_cfg,
p, fn, range,
map_in_fun{}, map_out_fun{},
std::move(xs)...));
}
actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::spawn_config& spawn_cfg,
const char* fn, const opencl::nd_range& range,
map_in_fun map_input, Ts&&... xs) const {
return actor_cast<actor>(impl::create(std::move(actor_cfg),
p, fn, spawn_cfg,
p, fn, range,
std::move(map_input),
map_out_fun{},
std::move(xs)...));
}
actor operator()(actor_config actor_cfg, const opencl::program_ptr p,
const char* fn, const opencl::spawn_config& spawn_cfg,
const char* fn, const opencl::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, spawn_cfg,
p, fn, range,
std::move(map_input),
std::move(map_output),
std::move(xs)...));
......
......@@ -113,9 +113,9 @@ public:
actor
>::type
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config, T x, Ts... xs) {
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, config,
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, range,
std::move(x), std::move(xs)...);
}
......@@ -130,10 +130,10 @@ public:
actor
>::type
spawn(const char* source, const char* fname,
const opencl::spawn_config& config, T x, Ts... xs) {
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, config,
create_program(source), fname, range,
std::move(x), std::move(xs)...);
}
......@@ -143,11 +143,11 @@ public:
/// `dims.empty()`, or `clCreateKernel` failed.
template <class Fun, class... Ts>
actor spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config,
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, config,
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, range,
std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
......@@ -159,12 +159,12 @@ public:
/// occured, or @p clCreateKernel failed.
template <class Fun, class... Ts>
actor spawn(const char* source, const char* fname,
const opencl::spawn_config& config,
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, config,
create_program(source), fname, range,
std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
......@@ -181,11 +181,11 @@ public:
actor
>::type
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config,
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, config,
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, range,
std::move(map_args), std::forward<T>(x), std::forward<Ts>(xs)...);
}
......@@ -200,12 +200,12 @@ public:
actor
>::type
spawn(const char* source, const char* fname,
const opencl::spawn_config& config,
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, config,
create_program(source), fname, range,
std::move(map_args), std::forward<T>(x), std::forward<Ts>(xs)...);
}
......@@ -222,11 +222,11 @@ public:
actor
>::type
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
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, config,
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, range,
std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
......@@ -242,12 +242,12 @@ public:
actor
>::type
spawn(const char* source, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
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, config,
create_program(source), fname, range,
std::move(map_args), std::move(map_result),
std::forward<Ts>(xs)...);
}
......@@ -262,11 +262,11 @@ public:
actor
>::type
spawn(const opencl::program_ptr prog, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
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, config,
return f(actor_config{system_.dummy_execution_unit()}, prog, fname, range,
std::move(map_args), std::forward<T>(x), std::forward<Ts>(xs)...);
}
......@@ -281,12 +281,12 @@ public:
actor
>::type
spawn(const char* source, const char* fname,
const opencl::spawn_config& config,
std::function<optional<message> (spawn_config&, message&)> map_args,
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, config,
create_program(source), fname, range,
std::move(map_args), std::forward<T>(x), std::forward<Ts>(xs)...);
}
......
......@@ -17,39 +17,39 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_SPAWN_CONFIG_HPP
#define CAF_OPENCL_SPAWN_CONFIG_HPP
#ifndef CAF_OPENCL_ND_RANGE_HPP
#define CAF_OPENCL_ND_RANGE_HPP
#include "caf/opencl/global.hpp"
namespace caf {
namespace opencl {
class spawn_config {
class nd_range {
public:
spawn_config(const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {})
: dims_{dims},
offset_{offset},
local_dims_{local_dims} {
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
}
spawn_config(opencl::dim_vec&& dims,
opencl::dim_vec&& offset = {},
opencl::dim_vec&& local_dims = {})
: dims_{std::move(dims)},
offset_{std::move(offset)},
local_dims_{std::move(local_dims)} {
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
}
spawn_config(const spawn_config&) = default;
spawn_config(spawn_config&&) = default;
nd_range(const nd_range&) = default;
nd_range(nd_range&&) = default;
spawn_config& operator=(const spawn_config&) = default;
spawn_config& operator=(spawn_config&&) = default;
nd_range& operator=(const nd_range&) = default;
nd_range& operator=(nd_range&&) = default;
const opencl::dim_vec& dimensions() const {
return dims_;
......@@ -72,4 +72,4 @@ private:
} // namespace opencl
} // namespace caf
#endif // CAF_OPENCL_SPAWN_CONFIG_HPP
#endif // CAF_OPENCL_ND_RANGE_HPP
......@@ -40,7 +40,7 @@
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/spawn_config.hpp"
#include "caf/opencl/nd_range.hpp"
namespace caf {
namespace opencl {
......@@ -101,7 +101,7 @@ public:
using input_types =
typename detail::tl_map<input_wrapped_types, extract_input_type>::type;
using input_mapping = typename std::conditional<PassConfig,
std::function<optional<message> (spawn_config&, message&)>,
std::function<optional<message> (nd_range&, message&)>,
std::function<optional<message> (message&)>
>::type;
......@@ -127,16 +127,16 @@ public:
}
static actor create(actor_config actor_conf, const program_ptr prog,
const char* kernel_name, const spawn_config& spawn_conf,
const char* kernel_name, const nd_range& range,
input_mapping map_args, output_mapping map_result,
Ts&&... xs) {
if (spawn_conf.dimensions().empty()) {
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() != spawn_conf.dimensions().size()) {
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";
......@@ -144,8 +144,8 @@ public:
throw std::runtime_error(oss.str());
}
};
check_vec(spawn_conf.offsets(), "offsets");
check_vec(spawn_conf.local_dimensions(), "local dimensions");
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()) {
......@@ -155,14 +155,14 @@ public:
false);
return make_actor<opencl_actor, actor>(sys.next_actor_id(), sys.node(),
&sys, std::move(actor_conf),
prog, kernel, spawn_conf,
prog, kernel, range,
std::move(map_args),
std::move(map_result),
std::forward_as_tuple(xs...));
}
return make_actor<opencl_actor, actor>(sys.next_actor_id(), sys.node(),
&sys, std::move(actor_conf),
prog, itr->second, spawn_conf,
prog, itr->second, range,
std::move(map_args),
std::move(map_result),
std::forward_as_tuple(xs...));
......@@ -203,7 +203,7 @@ public:
std::move(result_lengths),
std::move(content),
std::move(result),
config_
range_
);
cmd->enqueue();
}
......@@ -223,7 +223,7 @@ public:
}
opencl_actor(actor_config actor_conf, const program_ptr prog,
cl_kernel_ptr kernel, spawn_config spawn_conf,
cl_kernel_ptr kernel, nd_range range,
input_mapping map_args, output_mapping map_result,
std::tuple<Ts...> xs)
: monitorable_actor(actor_conf),
......@@ -231,13 +231,13 @@ public:
program_(prog->program_),
context_(prog->context_),
queue_(prog->queue_),
config_(std::move(spawn_conf)),
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(config_.dimensions()),
std::end(config_.dimensions()),
default_length_ = std::accumulate(std::begin(range_.dimensions()),
std::end(range_.dimensions()),
size_t{1},
std::multiplies<size_t>{});
}
......@@ -489,7 +489,7 @@ public:
template <bool Q = PassConfig>
typename std::enable_if<Q, bool>::type map_arguments(message& content) {
if (map_args_) {
auto mapped = map_args_(config_, content);
auto mapped = map_args_(range_, content);
if (!mapped) {
CAF_LOG_ERROR("Mapping argumentes failed.");
return false;
......@@ -503,7 +503,7 @@ public:
cl_program_ptr program_;
cl_context_ptr context_;
cl_command_queue_ptr queue_;
spawn_config config_;
nd_range range_;
input_mapping map_args_;
output_mapping map_results_;
std::tuple<Ts...> kernel_signature_;
......
......@@ -176,7 +176,7 @@ void multiplier(event_based_actor* self) {
// arguments for global memory to be contained in vectors,
// the vector is omitted here.
auto worker = mngr.spawn(kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
nd_range{dim_vec{matrix_size, matrix_size}},
unbox_args, box_res,
in<float>{}, in<float>{}, out<float>{});
......
......@@ -227,22 +227,22 @@ int main() {
return round_up((n + 1) / 2, half_block);
};
auto nd_conf = [half_block, get_size](size_t dim) {
return spawn_config{dim_vec{get_size(dim)}, {}, dim_vec{half_block}};
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 = spawn_config{dim_vec{half_block}, {}, dim_vec{half_block}};
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](spawn_config& conf, message& msg) -> optional<message> {
[nd_conf](nd_range& range, message& msg) -> optional<message> {
return msg.apply([&](uvec& vec) {
auto size = vec.size();
conf = nd_conf(size);
range = nd_conf(size);
return make_message(std::move(vec), static_cast<uval>(size));
});
},
......@@ -253,10 +253,10 @@ int main() {
);
auto phase2 = mngr.spawn(
prog, kernel_name_2, ndr,
[nd_conf](spawn_config& conf, message& msg) -> optional<message> {
[nd_conf](nd_range& range, message& msg) -> optional<message> {
return msg.apply([&](uref& data, uref& incs) {
auto size = incs.size();
conf = nd_conf(size);
range = nd_conf(size);
return make_message(move(data), move(incs), static_cast<uval>(size));
});
},
......@@ -266,10 +266,10 @@ int main() {
);
auto phase3 = mngr.spawn(
prog, kernel_name_3, ndr,
[nd_conf](spawn_config& conf, message& msg) -> optional<message> {
[nd_conf](nd_range& range, message& msg) -> optional<message> {
return msg.apply([&](uref& data, uref& incs) {
auto size = incs.size();
conf = nd_conf(size);
range = nd_conf(size);
return make_message(move(data), move(incs), static_cast<uval>(size));
});
},
......
......@@ -96,7 +96,7 @@ void multiplier(event_based_actor* self) {
// the vector type is omitted for brevity.
auto worker = self->system().opencl_manager().spawn(
kernel_source, kernel_name,
spawn_config{dim_vec{matrix_size, matrix_size}},
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
......
......@@ -341,7 +341,7 @@ void test_opencl(actor_system& sys) {
248, 286, 324, 362,
344, 398, 452, 506};
auto w1 = mngr.spawn(prog, kn_matrix,
opencl::spawn_config{dims{matrix_size, matrix_size}},
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 (
......@@ -351,9 +351,9 @@ void test_opencl(actor_system& sys) {
expected1, result);
}, others >> wrong_msg
);
opencl::spawn_config cfg2{dims{matrix_size, matrix_size}};
opencl::nd_range range2{dims{matrix_size, matrix_size}};
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, cfg2,
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 (
......@@ -374,9 +374,9 @@ void test_opencl(actor_system& sys) {
auto map_res = [](ivec result) -> message {
return make_message(matrix_type{move(result)});
};
opencl::spawn_config cfg3{dims{matrix_size, matrix_size}};
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, cfg3,
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>());
......@@ -387,8 +387,8 @@ void test_opencl(actor_system& sys) {
expected2.data(), result.data());
}, others >> wrong_msg
);
opencl::spawn_config cfg4{dims{matrix_size, matrix_size}};
auto w4 = mngr.spawn(prog, kn_matrix, cfg4,
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>());
......@@ -413,8 +413,8 @@ void test_opencl(actor_system& sys) {
// create program with opencl compiler flags
auto prog5 = mngr.create_program(kernel_source_compiler_flag, compiler_flag);
opencl::spawn_config cfg5{dims{array_size}};
auto w5 = mngr.spawn(prog5, kn_compiler_flag, cfg5,
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);
......@@ -434,13 +434,13 @@ void test_opencl(actor_system& sys) {
ivec arr6(reduce_buffer_size);
int n = static_cast<int>(arr6.capacity());
generate(arr6.begin(), arr6.end(), [&]{ return --n; });
opencl::spawn_config cfg6{dims{reduce_global_size},
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, cfg6,
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);
......@@ -460,7 +460,7 @@ void test_opencl(actor_system& sys) {
// constant memory arguments
const ivec arr7{problem_size};
auto w7 = mngr.spawn(kernel_source, kn_const,
opencl::spawn_config{dims{problem_size}},
opencl::nd_range{dims{problem_size}},
opencl::in<int>{},
opencl::out<int>{result_size_7});
self->send(w7, move(arr7));
......@@ -486,7 +486,7 @@ void test_arguments(actor_system& sys) {
const ivec expected1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto w1 = mngr.spawn(mngr.create_program(kernel_source, "", dev), kn_matrix,
opencl::spawn_config{dims{matrix_size, matrix_size}},
opencl::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 (
......@@ -498,7 +498,7 @@ void test_arguments(actor_system& sys) {
ivec expected9{input9};
for_each(begin(expected9), end(expected9), [](int& val){ val *= 2; });
auto w9 = mngr.spawn(kernel_source, kn_inout,
spawn_config{dims{problem_size}},
nd_range{dims{problem_size}},
opencl::in_out<int>{});
self->send(w9, move(input9));
self->receive(
......@@ -511,7 +511,7 @@ void test_arguments(actor_system& sys) {
for_each(begin(expected10), end(expected10), [](int& val){ val *= 2; });
auto result_size_10 = [=](const ivec& input) { return input.size(); };
auto w10 = mngr.spawn(kernel_source, kn_scratch,
spawn_config{dims{problem_size}},
nd_range{dims{problem_size}},
opencl::in_out<int>{},
opencl::scratch<int>{result_size_10});
self->send(w10, move(input10));
......@@ -533,7 +533,7 @@ void test_arguments(actor_system& sys) {
last += tmp;
}
auto work_local = mngr.spawn(kernel_source, kn_local,
spawn_config{dims{la_global}, {}, dims{la_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));
......@@ -545,7 +545,7 @@ void test_arguments(actor_system& sys) {
// Same test, different argument order
input_local = make_iota_vector<int>(la_global);
work_local = mngr.spawn(kernel_source, kn_order,
spawn_config{dims{la_global}, {}, dims{la_local}},
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));
......@@ -561,7 +561,7 @@ void test_arguments(actor_system& sys) {
for_each(begin(expected_private), end(expected_private),
[val_private](int& val){ val += val_private; });
auto worker_private = mngr.spawn(kernel_source, kn_private,
spawn_config{dims{problem_size}},
nd_range{dims{problem_size}},
opencl::in_out<int>{},
opencl::priv<int>{val_private});
self->send(worker_private, std::move(input_private));
......@@ -688,7 +688,7 @@ void test_in_val_out_val(actor_system& sys) {
};
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto 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) {
......@@ -723,8 +723,8 @@ void test_in_val_out_val(actor_system& sys) {
}, others >> wrong_msg);
// create program with opencl compiler flags
auto prog2 = mngr.create_program(kernel_source_compiler_flag, compiler_flag);
spawn_config conf2{dims{array_size}};
auto w4 = mngr.spawn(prog2, kn_compiler_flag, conf2,
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);
......@@ -742,9 +742,9 @@ void test_in_val_out_val(actor_system& sys) {
ivec input(reduce_buffer_size);
int n = static_cast<int>(input.capacity());
generate(input.begin(), input.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
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, conf3,
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);
......@@ -761,7 +761,7 @@ void test_in_val_out_val(actor_system& sys) {
// constant memory arguments
const ivec input2{problem_size};
auto w6 = mngr.spawn(kernel_source, kn_const,
spawn_config{dims{problem_size}},
nd_range{dims{problem_size}},
in<int>{}, out<int>{res_size2});
self->send(w6, move(input2));
ivec res5(problem_size);
......@@ -787,15 +787,15 @@ void test_in_val_out_mref(actor_system& sys) {
// tests
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, conf, in<int>{}, out<int, mref>{});
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, conf,
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) {
......@@ -812,9 +812,9 @@ void test_in_val_out_mref(actor_system& sys) {
ivec input(reduce_buffer_size);
int n = static_cast<int>(input.capacity());
generate(input.begin(), input.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
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, conf3,
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);
......@@ -842,8 +842,8 @@ void test_in_mref_out_val(actor_system& sys) {
// tests
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, conf, in<int, mref>{}, out<int>{});
auto 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);
......@@ -852,7 +852,7 @@ void test_in_mref_out_val(actor_system& sys) {
" (kernel wrapped in program)", res1, result);
}, others >> wrong_msg);
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, conf,
auto w2 = mngr.spawn(kernel_source, kn_matrix, range,
in<int, mref>{}, out<int, val>{});
self->send(w2, input1);
self->receive([&](const ivec& result) {
......@@ -869,9 +869,9 @@ void test_in_mref_out_val(actor_system& sys) {
ivec values(reduce_buffer_size);
int n = static_cast<int>(values.capacity());
generate(values.begin(), values.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
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, conf3,
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);
......@@ -900,8 +900,8 @@ void test_in_mref_out_mref(actor_system& sys) {
// tests
const ivec res1{ 56, 62, 68, 74, 152, 174, 196, 218,
248, 286, 324, 362, 344, 398, 452, 506};
auto conf = opencl::spawn_config{dims{matrix_size, matrix_size}};
auto w1 = mngr.spawn(prog, kn_matrix, conf,
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);
......@@ -911,7 +911,7 @@ void test_in_mref_out_mref(actor_system& sys) {
" (kernel wrapped in program)", res1, result);
}, others >> wrong_msg);
// Pass kernel directly to the actor
auto w2 = mngr.spawn(kernel_source, kn_matrix, conf,
auto w2 = mngr.spawn(kernel_source, kn_matrix, range,
in<int, mref>{}, out<int, mref>{});
self->send(w2, input1);
self->receive([&](iref& result) {
......@@ -928,9 +928,9 @@ void test_in_mref_out_mref(actor_system& sys) {
ivec values(reduce_buffer_size);
int n = static_cast<int>(values.capacity());
generate(values.begin(), values.end(), [&]{ return --n; });
spawn_config conf3{dims{reduce_global_size}, dims{}, dims{reduce_local_size}};
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, conf3,
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);
......@@ -959,17 +959,17 @@ void test_varying_arguments(actor_system& sys) {
};
// tests
size_t size = 23;
spawn_config conf{dims{size}};
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, conf,
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, conf,
auto w2 = mngr.spawn(prog, kn_varying, range,
in<int,mref>{}, out<int>{},
in<int>{}, out<int,mref>{});
self->send(w2, input2, input1);
......@@ -998,26 +998,26 @@ void test_inout(actor_system& sys) {
auto input3 = dev->global_argument(input);
ivec res{input};
for_each(begin(res), end(res), [](int& val){ val *= 2; });
auto conf = spawn_config{dims{problem_size}};
auto w1 = mngr.spawn(kernel_source, kn_inout, conf,
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, conf,
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, conf,
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, conf,
auto w4 = mngr.spawn(kernel_source, kn_inout, range,
in_out<int,mref,mref>{});
self->send(w4, input3);
self->receive([&](iref& result) {
......@@ -1039,18 +1039,18 @@ void test_priv(actor_system& sys) {
return sec::unexpected_message;
};
// tests
spawn_config conf{dims{problem_size}};
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, conf,
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, conf,
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) {
......@@ -1082,15 +1082,15 @@ void test_local(actor_system& sys) {
res[i] = last;
last += tmp;
}
auto conf = spawn_config{dims{global_size}, {}, dims{local_size}};
auto w = mngr.spawn(kernel_source, kn_local, conf,
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, conf,
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) {
......
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