Commit a0ae21f9 authored by Dominik Charousset's avatar Dominik Charousset

Merge branch 'topic/opencl' into unstable

parents 0f565fc5 e893f505
......@@ -233,7 +233,6 @@ if (ENABLE_OPENCL)
${LIBCPPA_SRC}
src/opencl/global.cpp
src/opencl/program.cpp
src/opencl/actor_facade.cpp
src/opencl/opencl_metainfo.cpp)
add_definitions(-DCPPA_OPENCL)
endif (ENABLE_OPENCL)
......
......@@ -56,6 +56,11 @@ namespace io {
class broker;
} // namespace io
namespace opencl {
template <typename Signature>
class actor_facade;
} // namespace opencl
namespace detail {
class raw_access;
} // namespace detail
......@@ -76,6 +81,11 @@ struct is_convertible_to_actor {
|| std::is_base_of<event_based_actor, T>::value;
};
template<typename T>
struct is_convertible_to_actor<opencl::actor_facade<T>> {
static constexpr bool value = true;
};
/**
* @brief Identifies an untyped actor.
*
......
......@@ -72,7 +72,7 @@ struct cl_spawn_helper<R (Ts...), void> {
using map_res_fun = typename impl::result_mapping;
template<typename... Us>
actor_ptr operator()(map_arg_fun f0,
actor operator()(map_arg_fun f0,
map_res_fun f1,
const opencl::program& p,
const char* fname,
......@@ -83,7 +83,7 @@ struct cl_spawn_helper<R (Ts...), void> {
}
template<typename... Us>
actor_ptr operator()(const opencl::program& p,
actor operator()(const opencl::program& p,
const char* fname,
Us&&... args) const {
using std::move;
......@@ -118,14 +118,15 @@ struct cl_spawn_helper<std::function<optional<cow_tuple<Ts...>> (any_tuple)>,
* failed.
*/
template<typename Signature, typename... Ts>
inline actor_ptr spawn_cl(const opencl::program& prog,
inline actor spawn_cl(const opencl::program& prog,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {}) {
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
using std::move;
detail::cl_spawn_helper<Signature> f;
return f(prog, fname, dims, offset, local_dims);
return f(prog, fname, dims, offset, local_dims, result_size);
}
/**
......@@ -136,17 +137,19 @@ inline actor_ptr spawn_cl(const opencl::program& prog,
* occured, or @p clCreateKernel failed.
*/
template<typename Signature, typename... Ts>
inline actor_ptr spawn_cl(const char* source,
inline actor spawn_cl(const char* source,
const char* fname,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {}) {
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
using std::move;
return spawn_cl<Signature, Ts...>(opencl::program::create(source),
fname,
dims,
offset,
local_dims);
local_dims,
result_size);
}
/**
......@@ -159,13 +162,14 @@ inline actor_ptr spawn_cl(const char* source,
* failed.
*/
template<typename MapArgs, typename MapResult>
inline actor_ptr spawn_cl(const opencl::program& prog,
inline actor spawn_cl(const opencl::program& prog,
const char* fname,
MapArgs map_args,
MapResult map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {}) {
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
using std::move;
typedef typename util::get_callable_trait<MapArgs>::fun_type f0;
typedef typename util::get_callable_trait<MapResult>::fun_type f1;
......@@ -176,7 +180,8 @@ inline actor_ptr spawn_cl(const opencl::program& prog,
fname,
dims,
offset,
local_dims);
local_dims,
result_size);
}
/**
......@@ -189,13 +194,14 @@ inline actor_ptr spawn_cl(const opencl::program& prog,
* occured, or @p clCreateKernel failed.
*/
template<typename MapArgs, typename MapResult>
inline actor_ptr spawn_cl(const char* source,
inline actor spawn_cl(const char* source,
const char* fun_name,
MapArgs map_args,
MapResult map_result,
const opencl::dim_vec& dims,
const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {}) {
const opencl::dim_vec& local_dims = {},
size_t result_size = 0) {
using std::move;
return spawn_cl(opencl::program::create(source),
fun_name,
......@@ -203,7 +209,8 @@ inline actor_ptr spawn_cl(const char* source,
move(map_result),
dims,
offset,
local_dims);
local_dims,
result_size);
}
} // namespace cppa
......
......@@ -32,9 +32,9 @@
#define CPPA_OPENCL_ACTOR_FACADE_HPP
#include <ostream>
#include <iostream>
#include <algorithm>
#include <stdexcept>
#include <iostream>
#include "cppa/cppa.hpp"
......@@ -52,8 +52,6 @@
#include "cppa/opencl/program.hpp"
#include "cppa/opencl/smart_ptr.hpp"
#include "cppa/detail/scheduled_actor_dummy.hpp"
namespace cppa {
namespace opencl {
......@@ -76,9 +74,10 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
typedef std::function<any_tuple(Ret&)> result_mapping;
static intrusive_ptr<actor_facade>
create(const program& prog, const char* kernel_name, arg_mapping map_args,
result_mapping map_result, const dim_vec& global_dims,
const dim_vec& offsets, const dim_vec& local_dims) {
create(const program& prog, const char* kernel_name,
arg_mapping map_args, result_mapping map_result,
const dim_vec& global_dims, const dim_vec& offsets,
const dim_vec& local_dims, size_t result_size) {
if (global_dims.empty()) {
auto str = "OpenCL kernel needs at least 1 global dimension.";
CPPA_LOGM_ERROR(detail::demangle(typeid(actor_facade)).c_str(),
......@@ -107,9 +106,15 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
oss.str());
throw std::runtime_error(oss.str());
}
return new actor_facade<Ret(Args...)>{
prog, kernel, global_dims, offsets,
local_dims, std::move(map_args), std::move(map_result)
if (result_size == 0) {
result_size = std::accumulate(global_dims.begin(),
global_dims.end(),
1,
std::multiplies<size_t>{});
}
return new actor_facade<Ret (Args...)>{
prog , kernel , global_dims , offsets,
local_dims, std::move(map_args), std::move(map_result), result_size
};
}
......@@ -121,41 +126,48 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
private:
using evnt_vec = std::vector<cl_event>;
using args_vec = std::vector<mem_ptr>;
actor_facade(const program& prog, kernel_ptr kernel,
const dim_vec& global_dimensions,
const dim_vec& global_offsets, const dim_vec& local_dimensions,
arg_mapping map_args, result_mapping map_result)
: m_kernel(kernel), m_program(prog.m_program),
m_context(prog.m_context), m_queue(prog.m_queue),
m_global_dimensions(global_dimensions),
m_global_offsets(global_offsets),
m_local_dimensions(local_dimensions), m_map_args(std::move(map_args)),
m_map_result(std::move(map_result)) {
const dim_vec& global_offsets,
const dim_vec& local_dimensions,
arg_mapping map_args, result_mapping map_result,
size_t result_size)
: m_kernel(kernel) , m_program(prog.m_program)
, m_context(prog.m_context) , m_queue(prog.m_queue)
, m_global_dimensions(global_dimensions)
, m_global_offsets(global_offsets)
, m_local_dimensions(local_dimensions)
, m_map_args(std::move(map_args))
, m_map_result(std::move(map_result))
, m_result_size(result_size)
{
CPPA_LOG_TRACE("id: " << this->id());
}
template <long... Is>
void enqueue_impl(const actor_ptr& sender, any_tuple msg, message_id id,
void enqueue_impl(const actor_addr& sender, any_tuple msg, message_id id,
util::int_list<Is...>) {
auto opt = m_map_args(std::move(msg));
if (opt) {
response_promise handle{ this, sender, id.response_id() };
size_t ret_size = std::accumulate(m_global_dimensions.begin(),
m_global_dimensions.end(), 1,
std::multiplies<size_t>{});
std::vector<mem_ptr> arguments;
add_arguments_to_kernel<Ret>(arguments, ret_size,
response_promise handle{this->address(), sender, id.response_id()};
evnt_vec events;
args_vec arguments;
add_arguments_to_kernel<Ret>(events, arguments, m_result_size,
get_ref<Is>(*opt)...);
auto cmd = make_counted<command<actor_facade, Ret>>(
handle, this, std::move(arguments));
handle, this,
std::move(events), std::move(arguments),
m_result_size, *opt
);
cmd->enqueue();
} else {
CPPA_LOGMF(CPPA_ERROR, this,
"actor_facade::enqueue() tuple_cast failed.");
CPPA_LOGMF(CPPA_ERROR, "actor_facade::enqueue() tuple_cast failed.");
}
}
typedef std::vector<mem_ptr> args_vec;
kernel_ptr m_kernel;
program_ptr m_program;
......@@ -166,56 +178,67 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
dim_vec m_local_dimensions;
arg_mapping m_map_args;
result_mapping m_map_result;
void add_arguments_to_kernel_rec(args_vec& arguments) {
cl_int err{ 0 };
for (size_t i = 1; i < arguments.size(); ++i) {
err = clSetKernelArg(m_kernel.get(), (i - 1), sizeof(cl_mem),
size_t m_result_size;
void add_arguments_to_kernel_rec(evnt_vec&, args_vec& arguments) {
cl_int err{0};
// rotate left (output buffer to the end)
rotate(begin(arguments), begin(arguments) + 1, end(arguments));
for(size_t i = 0; i < arguments.size(); ++i) {
err = clSetKernelArg(m_kernel.get(), i, sizeof(cl_mem),
static_cast<void*>(&arguments[i]));
CPPA_LOG_ERROR_IF(err != CL_SUCCESS,
"clSetKernelArg: " << get_opencl_error(err));
}
err = clSetKernelArg(m_kernel.get(), arguments.size() - 1,
sizeof(cl_mem), static_cast<void*>(&arguments[0]));
CPPA_LOG_ERROR_IF(err != CL_SUCCESS,
"clSetKernelArg: " << get_opencl_error(err));
clFlush(m_queue.get());
}
template <typename T0, typename... Ts>
void add_arguments_to_kernel_rec(args_vec& arguments, T0& arg0,
Ts&... args) {
cl_int err{ 0 };
auto buf = clCreateBuffer(
m_context.get(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(typename T0::value_type) * arg0.size(), arg0.data(), &err);
template<typename T0, typename... Ts>
void add_arguments_to_kernel_rec(evnt_vec& events, args_vec& arguments,
T0& arg0, Ts&... args) {
cl_int err{0};
size_t buffer_size = sizeof(typename T0::value_type) * arg0.size();
auto buffer = clCreateBuffer(m_context.get(),
CL_MEM_READ_ONLY,
buffer_size,
nullptr,
&err);
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, this,
"clCreateBuffer: " << get_opencl_error(err));
} else {
CPPA_LOGMF(CPPA_ERROR, "clCreateBuffer: " << get_opencl_error(err));
return;
}
cl_event event;
err = clEnqueueWriteBuffer(m_queue.get(), buffer, CL_FALSE, 0,
buffer_size, arg0.data(),
0, nullptr, &event);
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, "clEnqueueWriteBuffer: "
<< get_opencl_error(err));
return;
}
events.push_back(std::move(event));
mem_ptr tmp;
tmp.adopt(std::move(buf));
tmp.adopt(std::move(buffer));
arguments.push_back(tmp);
add_arguments_to_kernel_rec(arguments, args...);
}
add_arguments_to_kernel_rec(events, arguments, args...);
}
template <typename R, typename... Ts>
void add_arguments_to_kernel(args_vec& arguments, size_t ret_size,
Ts&&... args) {
template<typename R, typename... Ts>
void add_arguments_to_kernel(evnt_vec& events, args_vec& arguments,
size_t ret_size, Ts&&... args) {
arguments.clear();
cl_int err{ 0 };
auto buf = clCreateBuffer(m_context.get(), CL_MEM_WRITE_ONLY,
sizeof(typename R::value_type) * ret_size,
nullptr, &err);
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, this,
"clCreateBuffer: " << get_opencl_error(err));
} else {
CPPA_LOGMF(CPPA_ERROR, "clCreateBuffer: " << get_opencl_error(err));
return;
}
mem_ptr tmp;
tmp.adopt(std::move(buf));
arguments.push_back(tmp);
add_arguments_to_kernel_rec(arguments, std::forward<Ts>(args)...);
}
add_arguments_to_kernel_rec(events, arguments, std::forward<Ts>(args)...);
}
};
......
......@@ -53,39 +53,71 @@ class command : public ref_counted {
command(response_promise handle,
intrusive_ptr<T> actor_facade,
std::vector<mem_ptr> arguments)
: m_number_of_values(std::accumulate(actor_facade->m_global_dimensions.begin(),
actor_facade->m_global_dimensions.end(),
1, std::multiplies<size_t>{}))
std::vector<cl_event> events,
std::vector<mem_ptr> arguments,
size_t result_size,
any_tuple msg)
: m_result_size(result_size)
, m_handle(handle)
, m_actor_facade(actor_facade)
, m_queue(actor_facade->m_queue)
, m_arguments(move(arguments)) { }
, m_events(std::move(events))
, m_arguments(std::move(arguments))
, m_result(m_result_size)
, m_msg(msg) { }
~command() {
cl_int err{0};
for(auto& e : m_events) {
err = clReleaseEvent(e);
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, "clReleaseEvent: "
<< get_opencl_error(err));
}
}
}
void enqueue () {
CPPA_LOG_TRACE("command::enqueue()");
this->ref(); // reference held by the OpenCL comand queue
cl_int err{0};
auto event = m_kernel_event.get();
cl_event event_k;
auto data_or_nullptr = [](const dim_vec& vec) {
return vec.empty() ? nullptr : vec.data();
};
err = clEnqueueNDRangeKernel(m_queue.get(),
m_actor_facade->m_kernel.get(),
m_actor_facade->m_global_dimensions.size(),
data_or_nullptr(m_actor_facade->m_global_offsets),
data_or_nullptr(m_actor_facade->m_global_dimensions),
data_or_nullptr(m_actor_facade->m_local_dimensions),
0,
nullptr,
&event);
m_events.size(),
(m_events.empty() ? nullptr : m_events.data()),
&event_k);
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, self, "clEnqueueNDRangeKernel: "
CPPA_LOGMF(CPPA_ERROR, "clEnqueueNDRangeKernel: "
<< get_opencl_error(err));
this->deref(); // or can anything actually happen?
return;
}
err = clSetEventCallback(event,
else {
cl_event event_r;
err = clEnqueueReadBuffer(m_queue.get(),
m_arguments.back().get(),
CL_FALSE,
0,
sizeof(typename R::value_type) * m_result_size,
m_result.data(),
1,
&event_k,
&event_r);
if (err != CL_SUCCESS) {
throw std::runtime_error("clEnqueueReadBuffer: "
+ get_opencl_error(err));
this->deref(); // failed to enqueue command
return;
}
err = clSetEventCallback(event_r,
CL_COMPLETE,
[](cl_event, cl_int, void* data) {
auto cmd = reinterpret_cast<command*>(data);
......@@ -94,42 +126,34 @@ class command : public ref_counted {
},
this);
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, self, "clSetEventCallback: "
CPPA_LOGMF(CPPA_ERROR, "clSetEventCallback: "
<< get_opencl_error(err));
this->deref(); // callback is not set
return;
}
err = clFlush(m_queue.get());
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, self, "clFlush: " << get_opencl_error(err));
CPPA_LOGMF(CPPA_ERROR, "clFlush: " << get_opencl_error(err));
}
m_events.push_back(std::move(event_k));
m_events.push_back(std::move(event_r));
}
}
private:
int m_number_of_values;
int m_result_size;
response_promise m_handle;
intrusive_ptr<T> m_actor_facade;
event_ptr m_kernel_event;
command_queue_ptr m_queue;
std::vector<cl_event> m_events;
std::vector<mem_ptr> m_arguments;
R m_result;
any_tuple m_msg; // required to keep the argument buffers alive (for async copy)
void handle_results () {
cl_int err{0};
R result(m_number_of_values);
err = clEnqueueReadBuffer(m_queue.get(),
m_arguments[0].get(),
CL_TRUE,
0,
sizeof(typename R::value_type) * m_number_of_values,
result.data(),
0,
nullptr,
nullptr);
if (err != CL_SUCCESS) {
throw std::runtime_error("clEnqueueReadBuffer: "
+ get_opencl_error(err));
}
reply_tuple_to(m_handle, m_actor_facade->m_map_result(result));
m_handle.deliver(m_actor_facade->m_map_result(m_result));
}
};
......
......@@ -57,7 +57,7 @@ class program {
* from a given @p kernel_source.
* @returns A program object.
*/
static program create(const char* kernel_source, uint32_t device_id = 0);
static program create(const char* kernel_source, const char* options = nullptr, uint32_t device_id = 0);
private:
......
......@@ -142,7 +142,7 @@ inline bool operator!=(const square_matrix<Size>& lhs,
using matrix_type = square_matrix<matrix_size>;
void multiplier() {
void multiplier(event_based_actor* self) {
// create two matrices with ascending values
matrix_type m1;
......@@ -182,7 +182,7 @@ void multiplier() {
// send both matrices to the actor and
// wait for results in form of a matrix_type
sync_send(worker, move(m1), move(m2)).then(
self->sync_send(worker, move(m1), move(m2)).then(
[](const matrix_type& result) {
cout << "result:" << endl << to_string(result);
}
......@@ -194,7 +194,7 @@ int main() {
// it must be annouced to libcppa
announce<matrix_type>();
spawn(multiplier);
await_all_others_done();
await_all_actors_done();
shutdown();
return 0;
}
......@@ -74,7 +74,7 @@ void print_as_matrix(const vector<float>& matrix) {
}
}
void multiplier() {
void multiplier(event_based_actor* self) {
// the opencl actor only understands vectors
// so these vectors represent the matrices
vector<float> m1(matrix_size * matrix_size);
......@@ -99,11 +99,15 @@ void multiplier() {
// creates matrix_size * matrix_size global work items
// 4th arg: offsets for global dimensions (optional)
// 5th arg: local dimensions (optional)
// 6th arg: number of elements in the result buffer
auto worker = spawn_cl<float*(float*,float*)>(kernel_source,
kernel_name,
{matrix_size, matrix_size});
{matrix_size, matrix_size},
{},
{},
matrix_size * matrix_size);
// send both matrices to the actor and wait for a result
sync_send(worker, move(m1), move(m2)).then(
self->sync_send(worker, move(m1), move(m2)).then(
[](const vector<float>& result) {
cout << "result: " << endl;
print_as_matrix(result);
......@@ -114,7 +118,7 @@ void multiplier() {
int main() {
announce<vector<float>>();
spawn(multiplier);
await_all_others_done();
await_all_actors_done();
shutdown();
return 0;
}
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \L\ \/\ \__/\ \ \L\ \ \ \L\ \/\ \L\.\_ *
* /\____\\ \_\ \_,__/\ \____\\ \ ,__/\ \ ,__/\ \__/.\_\ *
* \/____/ \/_/\/___/ \/____/ \ \ \/ \ \ \/ \/__/\/_/ *
* \ \_\ \ \_\ *
* \/_/ \/_/ *
* *
* Copyright (C) 2011-2013 *
* Dominik Charousset <dominik.charousset@haw-hamburg.de> *
* Raphael Hiesgen <raphael.hiesgen@haw-hamburg.de> *
* *
* This file is part of libcppa. *
* libcppa is free software: you can redistribute it and/or modify it under *
* the terms of the GNU Lesser General Public License as published by the *
* Free Software Foundation; either version 2.1 of the License, *
* or (at your option) any later version. *
* *
* libcppa is distributed in the hope that it will be useful, *
* but WITHOUT ANY WARRANTY; without even the implied warranty of *
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. *
* See the GNU Lesser General Public License for more details. *
* *
* You should have received a copy of the GNU Lesser General Public License *
* along with libcppa. If not, see <http://www.gnu.org/licenses/>. *
\******************************************************************************/
#include "cppa/opencl/actor_facade.hpp"
#include "cppa/opencl/opencl_metainfo.hpp"
namespace cppa { namespace opencl {
} } // namespace cppa::opencl
......@@ -50,7 +50,7 @@ void opencl_metainfo::initialize()
ostringstream oss;
oss << "clGetPlatformIDs (getting number of platforms): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw logic_error(oss.str());
}
......@@ -62,7 +62,7 @@ void opencl_metainfo::initialize()
ostringstream oss;
oss << "clGetPlatformIDs (getting platform ids): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw logic_error(oss.str());
}
......@@ -81,7 +81,7 @@ void opencl_metainfo::initialize()
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetDeviceIDs: " << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str());
}
vector<cl_device_id> devices(num_devices);
......@@ -89,22 +89,30 @@ void opencl_metainfo::initialize()
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetDeviceIDs: " << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str());
}
auto pfn_notify = [](const char *errinfo,
const void *,
size_t,
void *) {
CPPA_LOG_ERROR("\n##### Error message via pfn_notify #####\n" +
string(errinfo) +
"\n########################################");
};
// create a context
m_context.adopt(clCreateContext(0,
devices.size(),
devices.data(),
nullptr,
pfn_notify,
nullptr,
&err));
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clCreateContext: " << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str());
}
......@@ -118,7 +126,7 @@ void opencl_metainfo::initialize()
char buf[buf_size];
err = clGetDeviceInfo(device.get(), CL_DEVICE_NAME, buf_size, buf, &return_size);
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, self, "clGetDeviceInfo (CL_DEVICE_NAME): "
CPPA_LOGMF(CPPA_ERROR, "clGetDeviceInfo (CL_DEVICE_NAME): "
<< get_opencl_error(err));
fill(buf, buf+buf_size, 0);
}
......@@ -128,7 +136,7 @@ void opencl_metainfo::initialize()
CL_QUEUE_PROFILING_ENABLE,
&err));
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_DEBUG, self, "Could not create command queue for device "
CPPA_LOGMF(CPPA_DEBUG, "Could not create command queue for device "
<< buf << ": " << get_opencl_error(err));
}
else {
......@@ -143,7 +151,7 @@ void opencl_metainfo::initialize()
oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_GROUP_SIZE): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str());
}
cl_uint max_work_item_dimensions = 0;
......@@ -157,7 +165,7 @@ void opencl_metainfo::initialize()
oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str());
}
dim_vec max_work_items_per_dim(max_work_item_dimensions);
......@@ -171,7 +179,7 @@ void opencl_metainfo::initialize()
oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_ITEM_SIZES): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str());
}
device_info dev_info{device,
......@@ -187,7 +195,7 @@ void opencl_metainfo::initialize()
ostringstream oss;
oss << "Could not create a command queue for "
<< "any present device.";
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str());
}
}
......
......@@ -45,7 +45,7 @@ namespace cppa { namespace opencl {
program::program(context_ptr context, command_queue_ptr queue, program_ptr program)
: m_context(move(context)), m_program(move(program)), m_queue(move(queue)) { }
program program::create(const char* kernel_source, uint32_t device_id) {
program program::create(const char* kernel_source, const char* options, uint32_t device_id) {
auto metainfo = get_opencl_metainfo();
auto devices = metainfo->get_devices();
auto context = metainfo->m_context;
......@@ -70,77 +70,21 @@ program program::create(const char* kernel_source, uint32_t device_id) {
&kernel_source,
&kernel_source_length,
&err));
if (err != CL_SUCCESS) {
throw runtime_error("clCreateProgramWithSource: "
+ get_opencl_error(err));
}
auto program_build_info = [&](size_t vs, void* v, size_t* vsr) {
return clGetProgramBuildInfo(pptr.get(),
devices[device_id].m_device.get(),
CL_PROGRAM_BUILD_LOG,
vs,
v,
vsr);
};
// build programm from program object
auto dev_tmp = devices[device_id].m_device.get();
err = clBuildProgram(pptr.get(), 1, &dev_tmp, nullptr, nullptr, nullptr);
err = clBuildProgram(pptr.get(), 1, &dev_tmp, options, nullptr, nullptr);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clBuildProgram: " << get_opencl_error(err);
size_t ret_size;
auto bi_err = program_build_info(0, nullptr, &ret_size);
if (bi_err == CL_SUCCESS) {
vector<char> build_log(ret_size);
bi_err = program_build_info(ret_size, build_log.data(), nullptr);
if (bi_err == CL_SUCCESS) {
if (ret_size <= 1) {
oss << " (no build log available)";
}
else {
build_log[ret_size - 1] = '\0';
cerr << build_log.data() << endl;
}
}
else { // program_build_info failed to get log
oss << " (with CL_PROGRAM_BUILD_LOG to get log)";
}
}
else { // program_build_info failed to get size of the log
oss << " (with CL_PROGRAM_BUILD_LOG to get size)";
}
CPPA_LOGM_ERROR(detail::demangle<program>().c_str(), oss.str());
// the build log will be printed by the
// pfn_notify (see opencl_metainfo.cpp)
throw runtime_error(oss.str());
}
# ifdef CPPA_DEBUG_MODE
else {
const char* where = "CL_PROGRAM_BUILD_LOG:get size";
size_t ret_size;
err = program_build_info(0, nullptr, &ret_size);
if (err == CL_SUCCESS) {
where = "CL_PROGRAM_BUILD_LOG:get log";
vector<char> build_log(ret_size+1);
err = program_build_info(ret_size, build_log.data(), nullptr);
if (err == CL_SUCCESS) {
if (ret_size > 1) {
CPPA_LOGM_ERROR(detail::demangle<program>().c_str(),
"clBuildProgram: error");
build_log[ret_size - 1] = '\0';
cerr << "clBuildProgram build log:\n"
<< build_log.data() << endl;
}
}
}
if (err != CL_SUCCESS) {
CPPA_LOGM_ERROR(detail::demangle<program>().c_str(),
"clGetProgramBuildInfo (" << where << "): "
<< get_opencl_error(err));
}
}
# endif
return {context, devices[device_id].m_cmd_queue, pptr};
}
......
......@@ -16,9 +16,28 @@ using namespace cppa::opencl;
namespace {
using ivec = vector<int>;
using fvec = vector<float>;
constexpr size_t matrix_size = 4;
constexpr size_t array_size = 32;
constexpr int magic_number = 23;
// since we do currently not support local memory arguments
// this size is fixed in the reduce kernel code
constexpr size_t reduce_buffer_size = 512 * 8;
constexpr size_t reduce_local_size = 512;
constexpr size_t reduce_work_groups = reduce_buffer_size / reduce_local_size;
constexpr size_t reduce_global_size = reduce_buffer_size;
constexpr size_t reduce_result_size = reduce_work_groups;
constexpr const char* kernel_name = "matrix_square";
constexpr const char* kernel_name_result_size = "result_size";
constexpr const char* kernel_name_compiler_flag = "compiler_flag";
constexpr const char* kernel_name_reduce = "reduce";
constexpr const char* kernel_name_const = "const_mod";
constexpr const char* compiler_flag = "-D OPENCL_CPPA_TEST_FLAG";
constexpr const char* kernel_source = R"__(
__kernel void matrix_square(__global int* matrix,
......@@ -35,11 +54,65 @@ constexpr const char* kernel_source = R"__(
)__";
constexpr const char* kernel_source_error = R"__(
__kernel void matrix_square(__global int*) {
__kernel void missing(__global int*) {
size_t semicolon
}
)__";
constexpr const char* kernel_source_compiler_flag = R"__(
__kernel void compiler_flag(__global int* input,
__global int* output) {
size_t x = get_global_id(0);
#ifdef OPENCL_CPPA_TEST_FLAG
output[x] = input[x];
#else
output[x] = 0;
#endif
}
)__";
// http://developer.amd.com/resources/documentation-articles/articles-whitepapers/
// opencl-optimization-case-study-simple-reductions
constexpr const char* kernel_source_reduce = R"__(
__kernel void reduce(__global int* buffer,
__global int* result) {
__local int scratch[512];
size_t length = get_global_size(0);
size_t global_index = get_global_id(0);
int accumulator = INFINITY;
// Loop sequentially over chunks of input vector
while (global_index < length) {
int element = buffer[global_index];
accumulator = (accumulator < element) ? accumulator : element;
global_index += length;
}
// Perform parallel reduction
int local_index = get_local_id(0);
scratch[local_index] = accumulator;
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];
}
}
)__";
constexpr const char* kernel_source_const = R"__(
__kernel void const_mod(__constant int* input,
__global int* output) {
size_t idx = get_global_id(0);
output[idx] = input[0];
}
)__";
}
template<size_t Size>
......@@ -102,11 +175,9 @@ inline bool operator!=(const square_matrix<Size>& lhs,
using matrix_type = square_matrix<matrix_size>;
int main() {
CPPA_TEST(test_opencl);
void test_opencl() {
announce<ivec>();
announce<matrix_type>();
scoped_actor self;
const ivec expected1{ 56, 62, 68, 74
, 152, 174, 196, 218
......@@ -117,9 +188,9 @@ int main() {
kernel_name,
{matrix_size, matrix_size});
ivec m1(matrix_size * matrix_size);
iota(m1.begin(), m1.end(), 0);
send(worker1, move(m1));
receive (
iota(begin(m1), end(m1), 0);
self->send(worker1, move(m1));
self->receive (
on_arg_match >> [&] (const ivec& result) {
CPPA_CHECK(equal(begin(expected1), end(expected1), begin(result)));
}
......@@ -129,9 +200,9 @@ int main() {
kernel_name,
{matrix_size, matrix_size});
ivec m2(matrix_size * matrix_size);
iota(m2.begin(), m2.end(), 0);
send(worker2, move(m2));
receive (
iota(begin(m2), end(m2), 0);
self->send(worker2, move(m2));
self->receive (
on_arg_match >> [&] (const ivec& result) {
CPPA_CHECK(equal(begin(expected1), end(expected1), begin(result)));
}
......@@ -157,8 +228,8 @@ int main() {
auto worker3 = spawn_cl(program::create(kernel_source), kernel_name,
map_args, map_results,
{matrix_size, matrix_size});
send(worker3, move(m3));
receive (
self->send(worker3, move(m3));
self->receive (
on_arg_match >> [&] (const matrix_type& result) {
CPPA_CHECK(expected2 == result);
}
......@@ -170,8 +241,8 @@ int main() {
map_args, map_results,
{matrix_size, matrix_size}
);
send(worker4, move(m4));
receive (
self->send(worker4, move(m4));
self->receive (
on_arg_match >> [&] (const matrix_type& result) {
CPPA_CHECK(expected2 == result);
}
......@@ -185,8 +256,65 @@ int main() {
CPPA_CHECK_EQUAL("clBuildProgram: CL_BUILD_PROGRAM_FAILURE", exc.what());
}
cppa::await_all_others_done();
cppa::shutdown();
// 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 worker5 = spawn_cl<ivec(ivec&)>(prog5, kernel_name_compiler_flag, {array_size});
self->send(worker5, move(arr5));
ivec expected3(array_size);
iota(begin(expected3), end(expected3), 0);
self->receive (
on_arg_match >> [&] (const ivec& result) {
CPPA_CHECK(equal(begin(expected3), end(expected3), begin(result)));
}
);
// test for manuel return size selection
ivec arr6(reduce_buffer_size);
int n{static_cast<int>(arr6.capacity())};
generate(begin(arr6), end(arr6), [&]{ return --n; });
auto worker6 = spawn_cl<ivec(ivec&)>(kernel_source_reduce,
kernel_name_reduce,
{reduce_global_size},
{},
{reduce_local_size},
reduce_result_size);
self->send(worker6, move(arr6));
fvec expected4{3584, 3072, 2560, 2048, 1536, 1024, 512, 0};
self->receive (
on_arg_match >> [&] (const ivec& result) {
CPPA_CHECK(equal(begin(expected4), end(expected4), begin(result)));
}
);
// constant memory arguments
ivec arr7{magic_number};
auto worker7 = spawn_cl<ivec(ivec&)>(kernel_source_const,
kernel_name_const,
{magic_number});
self->send(worker7, move(arr7));
ivec expected5(magic_number);
fill(begin(expected5), end(expected5), magic_number);
self->receive(
on_arg_match >> [&] (const ivec& result) {
CPPA_CHECK(equal(begin(expected5), end(expected5), begin(result)));
}
);
}
int main() {
CPPA_TEST(test_opencl);
announce<ivec>();
announce<matrix_type>();
test_opencl();
await_all_actors_done();
shutdown();
return CPPA_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