Commit c5e6d286 authored by Joseph Noir's avatar Joseph Noir

reduced overhead

added move to avoid copies, a device id can be chosen when a program
is created (all available devices are numbered from 0), added
device_info to store information about available devices
parent 676fc15c
......@@ -35,6 +35,7 @@
#include <ostream>
#include <algorithm>
#include <stdexcept>
#include <iostream>
#include "cppa/cppa.hpp"
......@@ -58,8 +59,6 @@ namespace cppa { namespace opencl {
class opencl_metainfo;
command_queue_ptr get_command_queue(uint32_t);
template<typename Signature>
class actor_facade;
......@@ -120,7 +119,7 @@ class actor_facade<Ret(Args...)> : public actor {
void enqueue(const message_header& hdr, any_tuple msg) override {
CPPA_LOG_TRACE("");
typename util::il_indices<util::type_list<Args...>>::type indices;
enqueue_impl(hdr.sender, msg, hdr.id, indices);
enqueue_impl(hdr.sender, std::move(msg), hdr.id, indices);
}
private:
......@@ -135,6 +134,7 @@ class actor_facade<Ret(Args...)> : public actor {
: 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)
......@@ -145,8 +145,11 @@ class actor_facade<Ret(Args...)> : public actor {
}
template<long... Is>
void enqueue_impl(const actor_ptr& sender, any_tuple msg, message_id id, util::int_list<Is...>) {
auto opt = m_map_args(msg);
void enqueue_impl(const actor_ptr& sender,
any_tuple msg,
message_id id,
util::int_list<Is...>) {
auto opt = m_map_args(std::move(msg));
if (opt) {
response_handle handle{this, sender, id.response_id()};
size_t ret_size = std::accumulate(m_global_dimensions.begin(),
......@@ -159,19 +162,10 @@ class actor_facade<Ret(Args...)> : public actor {
ret_size,
get_ref<Is>(*opt)...);
auto cmd = make_counted<command_impl<actor_facade, Ret>>(handle,
this,
std::move(arguments));
cmd->ref();
cl_command_queue cmd_q = get_command_queue(0).get(); // todo: get the id from program
cmd->enqueue(cmd_q);
clFlush(cmd_q);
// cl_int err{clFlush(cmd_q)};
// if (err != CL_SUCCESS) {
// ostringstream oss;
// oss << "clFlush: " << get_opencl_error(err);
// CPPA_LOGMF(CPPA_ERROR, self, oss.str());
// throw runtime_error(oss.str());
// }
this,
std::move(arguments),
m_queue);
cmd->enqueue();
}
else { CPPA_LOGMF(CPPA_ERROR, this, "actor_facade::enqueue() tuple_cast failed."); }
}
......@@ -181,6 +175,7 @@ class actor_facade<Ret(Args...)> : public actor {
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;
......
......@@ -52,7 +52,7 @@ class command : public ref_counted {
command* next;
virtual void enqueue(command_queue_ptr queue) = 0;
virtual void enqueue() = 0;
};
......@@ -60,7 +60,7 @@ class command_dummy : public command {
public:
void enqueue(command_queue_ptr) override { }
void enqueue() override { }
};
template<typename T, typename R>
......@@ -69,42 +69,41 @@ class command_impl : public command {
public:
command_impl(response_handle handle,
intrusive_ptr<T> af_ptr,
std::vector<mem_ptr> arguments)
: m_number_of_values(std::accumulate(af_ptr->m_global_dimensions.begin(),
af_ptr->m_global_dimensions.end(),
intrusive_ptr<T> actor_facade,
std::vector<mem_ptr> arguments,
command_queue_ptr queue)
: m_number_of_values(std::accumulate(actor_facade->m_global_dimensions.begin(),
actor_facade->m_global_dimensions.end(),
1, std::multiplies<size_t>{}))
, m_handle(handle)
, m_af_ptr(af_ptr)
, m_arguments(move(arguments))
{
}
, m_actor_facade(actor_facade)
, m_queue(queue)
, m_arguments(move(arguments)) { }
void enqueue (command_queue_ptr queue) override {
void enqueue () override {
CPPA_LOG_TRACE("command::enqueue()");
this->ref();
this->ref(); // reference held by the OpenCL comand queue
cl_int err{0};
m_queue = queue;
auto evnt = m_kernel_event.get();
auto event = m_kernel_event.get();
auto data_or_nullptr = [](const dim_vec& vec) {
return vec.empty() ? nullptr : vec.data();
};
/* enqueue kernel */
err = clEnqueueNDRangeKernel(m_queue.get(),
m_af_ptr->m_kernel.get(),
m_af_ptr->m_global_dimensions.size(),
data_or_nullptr(m_af_ptr->m_global_offsets),
data_or_nullptr(m_af_ptr->m_global_dimensions),
data_or_nullptr(m_af_ptr->m_local_dimensions),
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,
&evnt);
&event);
if (err != CL_SUCCESS) {
throw std::runtime_error("clEnqueueNDRangeKernel: "
+ get_opencl_error(err));
CPPA_LOGMF(CPPA_ERROR, self, "clEnqueueNDRangeKernel: "
<< get_opencl_error(err));
}
err = clSetEventCallback(evnt,
err = clSetEventCallback(event,
CL_COMPLETE,
[](cl_event, cl_int, void* data) {
auto cmd = reinterpret_cast<command_impl*>(data);
......@@ -113,8 +112,13 @@ class command_impl : public command {
},
this);
if (err != CL_SUCCESS) {
throw std::runtime_error("clSetEventCallback: "
+ get_opencl_error(err));
CPPA_LOGMF(CPPA_ERROR, self, "clSetEventCallback: "
<< get_opencl_error(err));
}
err = clFlush(m_queue.get());
if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, self, "clFlush: " << get_opencl_error(err));
}
}
......@@ -122,7 +126,7 @@ class command_impl : public command {
int m_number_of_values;
response_handle m_handle;
intrusive_ptr<T> m_af_ptr;
intrusive_ptr<T> m_actor_facade;
event_ptr m_kernel_event;
command_queue_ptr m_queue;
std::vector<mem_ptr> m_arguments;
......@@ -143,7 +147,7 @@ class command_impl : public command {
throw std::runtime_error("clEnqueueReadBuffer: "
+ get_opencl_error(err));
}
reply_tuple_to(m_handle, m_af_ptr->m_map_result(result));
reply_tuple_to(m_handle, m_actor_facade->m_map_result(result));
}
};
......
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \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 3 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/>. *
\******************************************************************************/
#ifndef DEVICE_INFO_HPP
#define DEVICE_INFO_HPP
#include "cppa/opencl/global.hpp"
#include "cppa/opencl/program.hpp"
#include "cppa/opencl/smart_ptr.hpp"
namespace cppa { namespace opencl {
class device_info {
friend class program;
public:
device_info(device_ptr device,
command_queue_ptr queue,
size_t work_group_size,
cl_uint dimensons,
const dim_vec& items_per_dimension)
: m_max_work_group_size(work_group_size)
, m_max_dimensions(dimensons)
, m_max_work_items_per_dim(items_per_dimension)
, m_device(device)
, m_cmd_queue(queue) { }
// todo: need getter, m_
// see actor.hpp
inline size_t get_max_work_group_size();
inline cl_uint get_max_dimensions();
inline dim_vec get_max_work_items_per_dim();
private:
size_t m_max_work_group_size;
cl_uint m_max_dimensions;
dim_vec m_max_work_items_per_dim;
device_ptr m_device;
command_queue_ptr m_cmd_queue;
};
/******************************************************************************\
* implementation of inline member functions *
\******************************************************************************/
inline size_t device_info::get_max_work_group_size() {
return m_max_work_group_size;
}
inline cl_uint device_info::get_max_dimensions() {
return m_max_dimensions;
}
inline dim_vec device_info::get_max_work_items_per_dim() {
return m_max_work_items_per_dim;
}
} } // namespace cppa::opencl
#endif // DEVICE_INFO_HPP
......@@ -42,6 +42,7 @@
#include "cppa/opencl/global.hpp"
#include "cppa/opencl/program.hpp"
#include "cppa/opencl/smart_ptr.hpp"
#include "cppa/opencl/device_info.hpp"
#include "cppa/opencl/actor_facade.hpp"
#include "cppa/detail/singleton_mixin.hpp"
......@@ -49,48 +50,6 @@
namespace cppa { namespace opencl {
//template<typename Ret, typename... Args>
//actor_ptr spawn(const program& prog,
// const char* kernel_name,
// const dim_vec& global_dims,
// const dim_vec& offsets,
// const dim_vec& local_dims,
// std::function<option<cow_tuple<typename util::rm_const_and_ref<Args>::type...>>(any_tuple)> map_args,
// std::function<any_tuple(Ret&)> map_result)
//{
// return actor_facade<Ret (Args...)>::create(prog,
// kernel_name,
// global_dims,
// offsets,
// local_dims,
// std::move(map_args),
// std::move(map_result));
//}
//template<typename Ret, typename... Args>
//actor_ptr spawn(const program& prog,
// const char* kernel_name,
// const dim_vec& global_dims,
// const dim_vec& offsets = {},
// const dim_vec& local_dims = {})
//{
// std::function<option<cow_tuple<typename util::rm_const_and_ref<Args>::type...>>(any_tuple)>
// map_args = [] (any_tuple msg) {
// return tuple_cast<typename util::rm_const_and_ref<Args>::type...>(msg);
// };
// std::function<any_tuple(Ret&)> map_result = [] (Ret& result) {
// return make_any_tuple(std::move(result));
// };
// return spawn<Ret, Args...>(prog,
// kernel_name,
// global_dims,
// offsets,
// local_dims,
// std::move(map_args),
// std::move(map_result));
//}
class opencl_metainfo {
friend class program;
......@@ -99,29 +58,9 @@ class opencl_metainfo {
public:
private:
const std::vector<device_info> get_devices() const;
struct device_info {
uint32_t id;
command_queue_ptr cmd_queue;
device_ptr dev_id;
size_t max_itms_per_grp;
cl_uint max_dim;
dim_vec max_itms_per_dim;
device_info(unsigned id,
command_queue_ptr queue,
device_ptr device_id,
size_t max_itms_per_grp,
cl_uint max_dim,
const dim_vec& max_itms_per_dim)
: id(id)
, cmd_queue(queue)
, dev_id(device_id)
, max_itms_per_grp(max_itms_per_grp)
, max_dim(max_dim)
, max_itms_per_dim(max_itms_per_dim) { }
};
private:
static inline opencl_metainfo* create_singleton() {
return new opencl_metainfo;
......@@ -131,8 +70,6 @@ class opencl_metainfo {
void dispose();
void destroy();
std::atomic<uint32_t> dev_id_gen;
context_ptr m_context;
std::vector<device_info> m_devices;
......
......@@ -61,11 +61,11 @@ class program {
private:
program(context_ptr context, program_ptr program, uint32_t device_id);
program(context_ptr context, command_queue_ptr queue, program_ptr program);
uint32_t m_device_id;
context_ptr m_context;
program_ptr m_program;
command_queue_ptr m_queue;
};
......
......@@ -34,9 +34,4 @@
namespace cppa { namespace opencl {
// todo: find device by id
command_queue_ptr get_command_queue(uint32_t) {
return get_opencl_metainfo()->m_devices.front().cmd_queue;
}
} } // namespace cppa::opencl
......@@ -34,6 +34,10 @@ using namespace std;
namespace cppa { namespace opencl {
const std::vector<device_info> opencl_metainfo::get_devices() const {
return m_devices;
}
void opencl_metainfo::initialize()
{
cl_int err{0};
......@@ -104,13 +108,13 @@ void opencl_metainfo::initialize()
CPPA_LOG_TRACE("Creating command queue for device(s).");
device_ptr device;
device.adopt(d);
uint32_t id{++dev_id_gen};
size_t return_size{0};
static constexpr size_t buf_size = 128;
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): " << get_opencl_error(err));
CPPA_LOGMF(CPPA_ERROR, self, "clGetDeviceInfo (CL_DEVICE_NAME): "
<< get_opencl_error(err));
fill(buf, buf+buf_size, 0);
}
command_queue_ptr cmd_queue;
......@@ -131,8 +135,8 @@ void opencl_metainfo::initialize()
&return_size);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetDeviceInfo (" << id
<< ":CL_DEVICE_MAX_WORK_GROUP_SIZE): "
oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_GROUP_SIZE): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
throw runtime_error(oss.str());
......@@ -145,8 +149,8 @@ void opencl_metainfo::initialize()
&return_size);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetDeviceInfo (" << id
<< ":CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): "
oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
throw runtime_error(oss.str());
......@@ -159,15 +163,14 @@ void opencl_metainfo::initialize()
&return_size);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "clGetDeviceInfo (" << id
<< ":CL_DEVICE_MAX_WORK_ITEM_SIZES): "
oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_ITEM_SIZES): "
<< get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str());
throw runtime_error(oss.str());
}
device_info dev_info{id,
device_info dev_info{device,
cmd_queue,
device,
max_work_group_size,
max_work_item_dimensions,
max_work_items_per_dim};
......@@ -194,7 +197,6 @@ void opencl_metainfo::dispose() {
opencl_metainfo* get_opencl_metainfo() {
return detail::singleton_manager::get_opencl_metainfo();
return nullptr;
}
} } // namespace cppa::opencl
......
......@@ -42,18 +42,30 @@ using namespace std;
namespace cppa { namespace opencl {
program::program(context_ptr context, program_ptr program, uint32_t device_id)
: m_device_id(device_id), m_context(move(context)), m_program(move(program)) { }
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) {
context_ptr cptr = get_opencl_metainfo()->m_context;
auto metainfo = get_opencl_metainfo();
auto devices = metainfo->get_devices();
auto context = metainfo->m_context;
if (devices.size() <= device_id) {
ostringstream oss;
oss << "Device id " << device_id
<< " is not a vaild device. Maximum id is: "
<< (devices.size() -1) << ".";
CPPA_LOGM_ERROR(detail::demangle<program>().c_str(), oss.str());
throw runtime_error(oss.str());
}
cl_int err{0};
// create program object from kernel source
size_t kernel_source_length = strlen(kernel_source);
program_ptr pptr;
pptr.adopt(clCreateProgramWithSource(cptr.get(),
pptr.adopt(clCreateProgramWithSource(context.get(),
1,
&kernel_source,
&kernel_source_length,
......@@ -77,7 +89,7 @@ program program::create(const char* kernel_source, uint32_t device_id) {
err = clBuildProgram(pptr.get(), 0, nullptr, nullptr, nullptr, nullptr);
if (err != CL_SUCCESS) {
// todo: chosoe device, not just front
device_ptr device{get_opencl_metainfo()->m_devices.front().dev_id};
device_ptr device{devices[device_id].m_device};
const char* where = "CL_PROGRAM_BUILD_LOG:get size";
size_t ret_size;
auto bi_err = program_build_info(device, 0, nullptr, &ret_size);
......@@ -108,7 +120,7 @@ program program::create(const char* kernel_source, uint32_t device_id) {
}
else {
# ifdef CPPA_DEBUG_MODE
device_ptr device{get_opencl_metainfo()->m_devices.front().dev_id};
device_ptr device{devices[device_id].m_device};
const char* where = "CL_PROGRAM_BUILD_LOG:get size";
size_t ret_size;
err = program_build_info(device, 0, nullptr, &ret_size);
......@@ -133,7 +145,7 @@ program program::create(const char* kernel_source, uint32_t device_id) {
}
# endif
}
return {cptr, pptr, device_id};
return {context, devices[device_id].m_cmd_queue, pptr};
}
} } // namespace cppa::opencl
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