Commit bb012cd3 authored by Dominik Charousset's avatar Dominik Charousset

Merge branch 'topic/opencl' into unstable

parents b95fe60d 2319ac39
...@@ -32,9 +32,9 @@ ...@@ -32,9 +32,9 @@
#define CPPA_OPENCL_ACTOR_FACADE_HPP #define CPPA_OPENCL_ACTOR_FACADE_HPP
#include <ostream> #include <ostream>
#include <iostream>
#include <algorithm> #include <algorithm>
#include <stdexcept> #include <stdexcept>
#include <iostream>
#include "cppa/cppa.hpp" #include "cppa/cppa.hpp"
...@@ -52,8 +52,6 @@ ...@@ -52,8 +52,6 @@
#include "cppa/opencl/program.hpp" #include "cppa/opencl/program.hpp"
#include "cppa/opencl/smart_ptr.hpp" #include "cppa/opencl/smart_ptr.hpp"
#include "cppa/detail/scheduled_actor_dummy.hpp"
namespace cppa { namespace cppa {
namespace opencl { namespace opencl {
...@@ -76,9 +74,10 @@ class actor_facade<Ret(Args...)> : public abstract_actor { ...@@ -76,9 +74,10 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
typedef std::function<any_tuple(Ret&)> result_mapping; typedef std::function<any_tuple(Ret&)> result_mapping;
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, arg_mapping map_args, result_mapping map_result,
const dim_vec& offsets, const dim_vec& local_dims) { const dim_vec& global_dims, const dim_vec& offsets,
const dim_vec& local_dims, size_t result_size) {
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.";
CPPA_LOGM_ERROR(detail::demangle(typeid(actor_facade)).c_str(), CPPA_LOGM_ERROR(detail::demangle(typeid(actor_facade)).c_str(),
...@@ -107,9 +106,15 @@ class actor_facade<Ret(Args...)> : public abstract_actor { ...@@ -107,9 +106,15 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
oss.str()); oss.str());
throw std::runtime_error(oss.str()); throw std::runtime_error(oss.str());
} }
return new actor_facade<Ret(Args...)>{ if (result_size == 0) {
prog, kernel, global_dims, offsets, result_size = std::accumulate(global_dims.begin(),
local_dims, std::move(map_args), std::move(map_result) 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 { ...@@ -121,41 +126,48 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
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_dimensions,
const dim_vec& global_offsets, const dim_vec& local_dimensions, const dim_vec& global_offsets,
arg_mapping map_args, result_mapping map_result) const dim_vec& local_dimensions,
: m_kernel(kernel), m_program(prog.m_program), arg_mapping map_args, result_mapping map_result,
m_context(prog.m_context), m_queue(prog.m_queue), size_t result_size)
m_global_dimensions(global_dimensions), : m_kernel(kernel) , m_program(prog.m_program)
m_global_offsets(global_offsets), , m_context(prog.m_context) , m_queue(prog.m_queue)
m_local_dimensions(local_dimensions), m_map_args(std::move(map_args)), , m_global_dimensions(global_dimensions)
m_map_result(std::move(map_result)) { , 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()); CPPA_LOG_TRACE("id: " << this->id());
} }
template <long... Is> 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...>) { util::int_list<Is...>) {
auto opt = m_map_args(std::move(msg)); auto opt = m_map_args(std::move(msg));
if (opt) { if (opt) {
response_promise handle{ this, sender, id.response_id() }; response_promise handle{this->address(), sender, id.response_id()};
size_t ret_size = std::accumulate(m_global_dimensions.begin(), evnt_vec events;
m_global_dimensions.end(), 1, args_vec arguments;
std::multiplies<size_t>{}); add_arguments_to_kernel<Ret>(events, arguments, m_result_size,
std::vector<mem_ptr> arguments;
add_arguments_to_kernel<Ret>(arguments, ret_size,
get_ref<Is>(*opt)...); get_ref<Is>(*opt)...);
auto cmd = make_counted<command<actor_facade, Ret>>( 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(); cmd->enqueue();
} else { } else {
CPPA_LOGMF(CPPA_ERROR, this, CPPA_LOGMF(CPPA_ERROR, "actor_facade::enqueue() tuple_cast failed.");
"actor_facade::enqueue() tuple_cast failed.");
} }
} }
typedef std::vector<mem_ptr> args_vec;
kernel_ptr m_kernel; kernel_ptr m_kernel;
program_ptr m_program; program_ptr m_program;
...@@ -166,56 +178,67 @@ class actor_facade<Ret(Args...)> : public abstract_actor { ...@@ -166,56 +178,67 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
dim_vec m_local_dimensions; dim_vec m_local_dimensions;
arg_mapping m_map_args; arg_mapping m_map_args;
result_mapping m_map_result; result_mapping m_map_result;
size_t m_result_size;
void add_arguments_to_kernel_rec(args_vec& arguments) {
cl_int err{ 0 }; void add_arguments_to_kernel_rec(evnt_vec&, args_vec& arguments) {
for (size_t i = 1; i < arguments.size(); ++i) { cl_int err{0};
err = clSetKernelArg(m_kernel.get(), (i - 1), sizeof(cl_mem), // 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])); static_cast<void*>(&arguments[i]));
CPPA_LOG_ERROR_IF(err != CL_SUCCESS, CPPA_LOG_ERROR_IF(err != CL_SUCCESS,
"clSetKernelArg: " << get_opencl_error(err)); "clSetKernelArg: " << get_opencl_error(err));
} }
err = clSetKernelArg(m_kernel.get(), arguments.size() - 1, clFlush(m_queue.get());
sizeof(cl_mem), static_cast<void*>(&arguments[0]));
CPPA_LOG_ERROR_IF(err != CL_SUCCESS,
"clSetKernelArg: " << get_opencl_error(err));
} }
template <typename T0, typename... Ts> template<typename T0, typename... Ts>
void add_arguments_to_kernel_rec(args_vec& arguments, T0& arg0, void add_arguments_to_kernel_rec(evnt_vec& events, args_vec& arguments,
Ts&... args) { T0& arg0, Ts&... args) {
cl_int err{ 0 }; cl_int err{0};
auto buf = clCreateBuffer( size_t buffer_size = sizeof(typename T0::value_type) * arg0.size();
m_context.get(), CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, auto buffer = clCreateBuffer(m_context.get(),
sizeof(typename T0::value_type) * arg0.size(), arg0.data(), &err); CL_MEM_READ_ONLY,
buffer_size,
nullptr,
&err);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, this, CPPA_LOGMF(CPPA_ERROR, "clCreateBuffer: " << get_opencl_error(err));
"clCreateBuffer: " << get_opencl_error(err)); return;
} else {
mem_ptr tmp;
tmp.adopt(std::move(buf));
arguments.push_back(tmp);
add_arguments_to_kernel_rec(arguments, args...);
} }
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(buffer));
arguments.push_back(tmp);
add_arguments_to_kernel_rec(events, arguments, args...);
} }
template <typename R, typename... Ts> template<typename R, typename... Ts>
void add_arguments_to_kernel(args_vec& arguments, size_t ret_size, void add_arguments_to_kernel(evnt_vec& events, args_vec& arguments,
Ts&&... args) { size_t ret_size, Ts&&... args) {
arguments.clear(); arguments.clear();
cl_int err{ 0 }; cl_int err{ 0 };
auto buf = clCreateBuffer(m_context.get(), CL_MEM_WRITE_ONLY, auto buf = clCreateBuffer(m_context.get(), CL_MEM_WRITE_ONLY,
sizeof(typename R::value_type) * ret_size, sizeof(typename R::value_type) * ret_size,
nullptr, &err); nullptr, &err);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, this, CPPA_LOGMF(CPPA_ERROR, "clCreateBuffer: " << get_opencl_error(err));
"clCreateBuffer: " << get_opencl_error(err)); return;
} else {
mem_ptr tmp;
tmp.adopt(std::move(buf));
arguments.push_back(tmp);
add_arguments_to_kernel_rec(arguments, std::forward<Ts>(args)...);
} }
mem_ptr tmp;
tmp.adopt(std::move(buf));
arguments.push_back(tmp);
add_arguments_to_kernel_rec(events, arguments, std::forward<Ts>(args)...);
} }
}; };
......
...@@ -53,83 +53,107 @@ class command : public ref_counted { ...@@ -53,83 +53,107 @@ class command : public ref_counted {
command(response_promise handle, command(response_promise handle,
intrusive_ptr<T> actor_facade, intrusive_ptr<T> actor_facade,
std::vector<mem_ptr> arguments) std::vector<cl_event> events,
: m_number_of_values(std::accumulate(actor_facade->m_global_dimensions.begin(), std::vector<mem_ptr> arguments,
actor_facade->m_global_dimensions.end(), size_t result_size,
1, std::multiplies<size_t>{})) any_tuple msg)
: m_result_size(result_size)
, m_handle(handle) , m_handle(handle)
, m_actor_facade(actor_facade) , m_actor_facade(actor_facade)
, m_queue(actor_facade->m_queue) , 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 () { void enqueue () {
CPPA_LOG_TRACE("command::enqueue()"); CPPA_LOG_TRACE("command::enqueue()");
this->ref(); // reference held by the OpenCL comand queue this->ref(); // reference held by the OpenCL comand queue
cl_int err{0}; cl_int err{0};
auto event = m_kernel_event.get(); cl_event event_k;
auto data_or_nullptr = [](const dim_vec& vec) { auto data_or_nullptr = [](const dim_vec& vec) {
return vec.empty() ? nullptr : vec.data(); return vec.empty() ? nullptr : vec.data();
}; };
err = clEnqueueNDRangeKernel(m_queue.get(), err = clEnqueueNDRangeKernel(m_queue.get(),
m_actor_facade->m_kernel.get(), m_actor_facade->m_kernel.get(),
m_actor_facade->m_global_dimensions.size(), m_actor_facade->m_global_dimensions.size(),
data_or_nullptr(m_actor_facade->m_global_offsets), 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_global_dimensions),
data_or_nullptr(m_actor_facade->m_local_dimensions), data_or_nullptr(m_actor_facade->m_local_dimensions),
0, m_events.size(),
nullptr, (m_events.empty() ? nullptr : m_events.data()),
&event); &event_k);
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
CPPA_LOGMF(CPPA_ERROR, self, "clEnqueueNDRangeKernel: " CPPA_LOGMF(CPPA_ERROR, "clEnqueueNDRangeKernel: "
<< get_opencl_error(err)); << get_opencl_error(err));
this->deref(); // or can anything actually happen?
return;
} }
else {
err = clSetEventCallback(event, cl_event event_r;
CL_COMPLETE, err = clEnqueueReadBuffer(m_queue.get(),
[](cl_event, cl_int, void* data) { m_arguments.back().get(),
auto cmd = reinterpret_cast<command*>(data); CL_FALSE,
cmd->handle_results(); 0,
cmd->deref(); sizeof(typename R::value_type) * m_result_size,
}, m_result.data(),
this); 1,
if (err != CL_SUCCESS) { &event_k,
CPPA_LOGMF(CPPA_ERROR, self, "clSetEventCallback: " &event_r);
<< get_opencl_error(err)); if (err != CL_SUCCESS) {
} throw std::runtime_error("clEnqueueReadBuffer: "
+ get_opencl_error(err));
err = clFlush(m_queue.get()); this->deref(); // failed to enqueue command
if (err != CL_SUCCESS) { return;
CPPA_LOGMF(CPPA_ERROR, self, "clFlush: " << get_opencl_error(err)); }
err = clSetEventCallback(event_r,
CL_COMPLETE,
[](cl_event, cl_int, void* data) {
auto cmd = reinterpret_cast<command*>(data);
cmd->handle_results();
cmd->deref();
},
this);
if (err != CL_SUCCESS) {
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, "clFlush: " << get_opencl_error(err));
}
m_events.push_back(std::move(event_k));
m_events.push_back(std::move(event_r));
} }
} }
private: private:
int m_number_of_values; int m_result_size;
response_promise m_handle; response_promise m_handle;
intrusive_ptr<T> m_actor_facade; intrusive_ptr<T> m_actor_facade;
event_ptr m_kernel_event;
command_queue_ptr m_queue; command_queue_ptr m_queue;
std::vector<cl_event> m_events;
std::vector<mem_ptr> m_arguments; 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 () { void handle_results () {
cl_int err{0}; m_handle.deliver(m_actor_facade->m_map_result(m_result));
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));
} }
}; };
......
...@@ -57,7 +57,7 @@ class program { ...@@ -57,7 +57,7 @@ class program {
* from a given @p kernel_source. * from a given @p kernel_source.
* @returns A program object. * @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: private:
......
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \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() ...@@ -50,7 +50,7 @@ void opencl_metainfo::initialize()
ostringstream oss; ostringstream oss;
oss << "clGetPlatformIDs (getting number of platforms): " oss << "clGetPlatformIDs (getting number of platforms): "
<< get_opencl_error(err); << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw logic_error(oss.str()); throw logic_error(oss.str());
} }
...@@ -62,7 +62,7 @@ void opencl_metainfo::initialize() ...@@ -62,7 +62,7 @@ void opencl_metainfo::initialize()
ostringstream oss; ostringstream oss;
oss << "clGetPlatformIDs (getting platform ids): " oss << "clGetPlatformIDs (getting platform ids): "
<< get_opencl_error(err); << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw logic_error(oss.str()); throw logic_error(oss.str());
} }
...@@ -81,7 +81,7 @@ void opencl_metainfo::initialize() ...@@ -81,7 +81,7 @@ void opencl_metainfo::initialize()
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
ostringstream oss; ostringstream oss;
oss << "clGetDeviceIDs: " << get_opencl_error(err); oss << "clGetDeviceIDs: " << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
vector<cl_device_id> devices(num_devices); vector<cl_device_id> devices(num_devices);
...@@ -89,22 +89,30 @@ void opencl_metainfo::initialize() ...@@ -89,22 +89,30 @@ void opencl_metainfo::initialize()
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
ostringstream oss; ostringstream oss;
oss << "clGetDeviceIDs: " << get_opencl_error(err); oss << "clGetDeviceIDs: " << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_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 // create a context
m_context.adopt(clCreateContext(0, m_context.adopt(clCreateContext(0,
devices.size(), devices.size(),
devices.data(), devices.data(),
nullptr, pfn_notify,
nullptr, nullptr,
&err)); &err));
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
ostringstream oss; ostringstream oss;
oss << "clCreateContext: " << get_opencl_error(err); oss << "clCreateContext: " << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
...@@ -118,8 +126,8 @@ void opencl_metainfo::initialize() ...@@ -118,8 +126,8 @@ void opencl_metainfo::initialize()
char buf[buf_size]; char buf[buf_size];
err = clGetDeviceInfo(device.get(), CL_DEVICE_NAME, buf_size, buf, &return_size); err = clGetDeviceInfo(device.get(), CL_DEVICE_NAME, buf_size, buf, &return_size);
if (err != CL_SUCCESS) { 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)); << get_opencl_error(err));
fill(buf, buf+buf_size, 0); fill(buf, buf+buf_size, 0);
} }
command_queue_ptr cmd_queue; command_queue_ptr cmd_queue;
...@@ -128,8 +136,8 @@ void opencl_metainfo::initialize() ...@@ -128,8 +136,8 @@ void opencl_metainfo::initialize()
CL_QUEUE_PROFILING_ENABLE, CL_QUEUE_PROFILING_ENABLE,
&err)); &err));
if (err != CL_SUCCESS) { 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)); << buf << ": " << get_opencl_error(err));
} }
else { else {
size_t max_work_group_size{0}; size_t max_work_group_size{0};
...@@ -143,7 +151,7 @@ void opencl_metainfo::initialize() ...@@ -143,7 +151,7 @@ void opencl_metainfo::initialize()
oss << "clGetDeviceInfo (" oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_GROUP_SIZE): " << "CL_DEVICE_MAX_WORK_GROUP_SIZE): "
<< get_opencl_error(err); << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
cl_uint max_work_item_dimensions = 0; cl_uint max_work_item_dimensions = 0;
...@@ -157,7 +165,7 @@ void opencl_metainfo::initialize() ...@@ -157,7 +165,7 @@ void opencl_metainfo::initialize()
oss << "clGetDeviceInfo (" oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): " << "CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): "
<< get_opencl_error(err); << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
dim_vec max_work_items_per_dim(max_work_item_dimensions); dim_vec max_work_items_per_dim(max_work_item_dimensions);
...@@ -171,7 +179,7 @@ void opencl_metainfo::initialize() ...@@ -171,7 +179,7 @@ void opencl_metainfo::initialize()
oss << "clGetDeviceInfo (" oss << "clGetDeviceInfo ("
<< "CL_DEVICE_MAX_WORK_ITEM_SIZES): " << "CL_DEVICE_MAX_WORK_ITEM_SIZES): "
<< get_opencl_error(err); << get_opencl_error(err);
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
device_info dev_info{device, device_info dev_info{device,
...@@ -187,7 +195,7 @@ void opencl_metainfo::initialize() ...@@ -187,7 +195,7 @@ void opencl_metainfo::initialize()
ostringstream oss; ostringstream oss;
oss << "Could not create a command queue for " oss << "Could not create a command queue for "
<< "any present device."; << "any present device.";
CPPA_LOGMF(CPPA_ERROR, self, oss.str()); CPPA_LOGMF(CPPA_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
} }
......
...@@ -45,7 +45,7 @@ namespace cppa { namespace opencl { ...@@ -45,7 +45,7 @@ namespace cppa { namespace opencl {
program::program(context_ptr context, command_queue_ptr queue, program_ptr 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)) { } : 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 metainfo = get_opencl_metainfo();
auto devices = metainfo->get_devices(); auto devices = metainfo->get_devices();
auto context = metainfo->m_context; auto context = metainfo->m_context;
...@@ -70,77 +70,21 @@ program program::create(const char* kernel_source, uint32_t device_id) { ...@@ -70,77 +70,21 @@ program program::create(const char* kernel_source, uint32_t device_id) {
&kernel_source, &kernel_source,
&kernel_source_length, &kernel_source_length,
&err)); &err));
if (err != CL_SUCCESS) { if (err != CL_SUCCESS) {
throw runtime_error("clCreateProgramWithSource: " throw runtime_error("clCreateProgramWithSource: "
+ get_opencl_error(err)); + 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 // build programm from program object
auto dev_tmp = devices[device_id].m_device.get(); 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) { if (err != CL_SUCCESS) {
ostringstream oss; ostringstream oss;
oss << "clBuildProgram: " << get_opencl_error(err); oss << "clBuildProgram: " << get_opencl_error(err);
size_t ret_size; // the build log will be printed by the
auto bi_err = program_build_info(0, nullptr, &ret_size); // pfn_notify (see opencl_metainfo.cpp)
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());
throw runtime_error(oss.str()); 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}; return {context, devices[device_id].m_cmd_queue, pptr};
} }
......
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