Commit 0daa2a3f authored by Joseph Noir's avatar Joseph Noir

basic functionality added to cppa/opencl

parent a22482d3
......@@ -32,14 +32,272 @@
#ifndef CPPA_OPENCL_ACTOR_FACADE_HPP
#define CPPA_OPENCL_ACTOR_FACADE_HPP
#include <iostream>
#include "cppa/to_string.hpp"
#include "cppa/channel.hpp"
#include "cppa/tuple_cast.hpp"
#include "cppa/opencl/global.hpp"
#include "cppa/opencl/command.hpp"
#include "cppa/opencl/program.hpp"
#include "cppa/util/apply_args.hpp"
#include "cppa/opencl/smart_ptr.hpp"
#include "cppa/detail/scheduled_actor_dummy.hpp"
namespace cppa { namespace opencl {
class actor_facade {
class command_dispatcher;
void enqueue_to_dispatcher(command_dispatcher*, command*);
template<typename Signature>
class actor_facade;
template<typename Ret, typename... Args>
class actor_facade<Ret(Args...)> : public cppa::detail::scheduled_actor_dummy {
public:
actor_facade(command_dispatcher* dispatcher,
const program& prog,
const std::string& kernel_name)
: m_program(prog.m_program)
, m_kernel()
, m_context(prog.m_context)
, m_kernel_name(kernel_name)
, m_dispatcher(dispatcher){
}
void enqueue(actor* sender, any_tuple msg) {
auto opt = tuple_cast<Args...>(msg);
if (opt) {
response_handle handle{this, sender, message_id_t{}};
auto fun = [=](cl_command_queue arg) {
auto indices = util::get_indices(*opt);
return util::apply_args_prefixed(*this,
*opt,
indices,
arg,
handle,
1024,
1,
1,
0,
1,
1);
};
enqueue_to_dispatcher(m_dispatcher, new command(fun, sender));
}
else {
// slap caller around with a large fish
}
}
/*
* todo:
* local_work_group_size currently unused, passing nullptr instead
*/
void operator()(cl_command_queue queue,
response_handle handle,
size_t global_work_size_dim1,
size_t global_work_size_dim2,
size_t global_work_size_dim3,
size_t local_work_size_dim1,
size_t local_work_size_dim2,
size_t local_work_size_dim3,
Args... args) {
cl_int err{0};
if (!m_kernel.get()) {
std::cout << "no kernel found" << std::endl;
m_kernel.adopt(clCreateKernel(m_program.get(),
m_kernel_name.c_str(),
&err));
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clCreateKernel: '"
+ get_opencl_error(err)
+ "'.");
}
}
else {
std::cout << "i found a kernel" << std::endl;
}
cl_event event;
int number_of_values = global_work_size_dim1
* global_work_size_dim2
* global_work_size_dim3;
/* add arguments */
Ret result_buf(number_of_values);
std::vector<mem_ptr> arguments;
add_arguments_to_kernel(arguments,
m_context.get(),
m_kernel.get(),
queue,
result_buf,
std::forward<Args>(args)...);
/* enqueue kernel */
size_t global_work_group_size[] = {global_work_size_dim1,
global_work_size_dim2,
global_work_size_dim3};
if (local_work_size_dim1 == 0) {
err = clEnqueueNDRangeKernel(queue,
m_kernel.get(),
3,
NULL,
global_work_group_size,
nullptr,
0,
nullptr,
&event);
}
else {
size_t local_work_group_size[] = {local_work_size_dim1,
local_work_size_dim2,
local_work_size_dim3};
err = clEnqueueNDRangeKernel(queue,
m_kernel.get(),
3,
NULL,
global_work_group_size,
local_work_group_size,
0,
nullptr,
&event);
}
clReleaseEvent(event);
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clEnqueueNDRangeKernel: '"
+ get_opencl_error(err)
+ "'.");
}
clFinish(queue);
/* get results from gpu */
Ret results(number_of_values);
err = clEnqueueReadBuffer(queue,
arguments[0].get(),
CL_TRUE,
0,
sizeof(typename Ret::value_type) * number_of_values,
results.data(),
0,
NULL,
&event);
clReleaseEvent(event);
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clEnqueueReadBuffer: '"
+ get_opencl_error(err)
+ "'.");
}
reply_to(handle, results);
//return results;
}
private:
typedef std::vector<mem_ptr> args_vec;
program_ptr m_program;
kernel_ptr m_kernel;
context_ptr m_context;
std::string m_kernel_name;
command_dispatcher* m_dispatcher;
static void add_arguments_to_kernel_rec(args_vec& arguments,
cl_context,
cl_kernel kernel,
cl_command_queue queue) {
cl_int err{0};
for(unsigned long i{1}; i < arguments.size(); ++i) {
err = clSetKernelArg(kernel,
(i-1),
sizeof(cl_mem),
static_cast<void*>(&arguments[i]));
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clSetKernelArg: '"
+ get_opencl_error(err)
+ "'.");
}
}
err = clSetKernelArg(kernel,
arguments.size()-1,
sizeof(cl_mem),
static_cast<void*>(&arguments[0]));
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clSetKernelArg: '"
+ get_opencl_error(err)
+ "'.");
}
clFinish(queue);
}
template<typename T0, typename... Ts>
static void add_arguments_to_kernel_rec(args_vec& arguments,
cl_context context,
cl_kernel kernel,
cl_command_queue queue,
T0&& arg0,
Ts&&... args) {
cl_int err{0};
auto buf = clCreateBuffer(context,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
sizeof(typename T0::value_type)*arg0.size(),
arg0.data(),
&err);
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clCreateBuffer: '"
+ get_opencl_error(err)
+ "'.");
}
else {
mem_ptr tmp;
tmp.adopt(std::move(buf));
arguments.push_back(tmp);
return add_arguments_to_kernel_rec(arguments,
context,
kernel,
queue,
std::forward<Ts>(args)...);
}
}
template<typename R, typename... Ts>
static void add_arguments_to_kernel(args_vec& arguments,
cl_context context,
cl_kernel kernel,
cl_command_queue queue,
R& ret,
Ts&&... args) {
arguments.clear();
cl_int err{0};
auto buf = clCreateBuffer(context,
CL_MEM_WRITE_ONLY,
sizeof(typename R::value_type)*ret.size(),
nullptr,
&err);
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clCreateBuffer: '"
+ get_opencl_error(err)
+ "'.");
}
else {
mem_ptr tmp;
tmp.adopt(std::move(buf));
arguments.push_back(tmp);
return add_arguments_to_kernel_rec(arguments,
context,
kernel,
queue,
std::forward<Ts>(args)...);
}
}
};
} // namespace cppa::opencl
} } // namespace cppa::opencl
#endif // CPPA_OPENCL_ACTOR_FACADE_HPP
......@@ -32,14 +32,38 @@
#ifndef CPPA_OPENCL_COMMAND_HPP
#define CPPA_OPENCL_COMMAND_HPP
#include <vector>
#include <functional>
#include "cppa/actor.hpp"
#include "cppa/opencl/global.hpp"
namespace cppa { namespace opencl {
class command {
public:
command* next;
// bool is_new_job; // if false == call handle_cl_result
// struct callbacks {
// std::function<cl_job_id(cl_command_queue)> enqueue_cl_task;
// std::function<void()> handle_cl_result;
// };
// union { callbacks cbs; cl_job_id jid; };
std::function<void(cl_command_queue)> fun;
actor* sender;
template<typename T>
inline command(T f, actor* sender)
: next(nullptr), fun(f), sender(sender) { }
inline command() : next(nullptr), sender(nullptr) { }
};
} // namespace cppa::opencl
} } // namespace cppa::opencl
#endif // CPPA_OPENCL_COMMAND_HPP
......@@ -32,23 +32,98 @@
#ifndef CPPA_OPENCL_COMMAND_DISPATCHER_HPP
#define CPPA_OPENCL_COMMAND_DISPATCHER_HPP
#include <atomic>
#include <vector>
#include <functional>
#include "cppa/channel.hpp"
#include "cppa/opencl/global.hpp"
#include "cppa/opencl/command.hpp"
#include "cppa/opencl/program.hpp"
#include "cppa/opencl/smart_ptr.hpp"
#include "cppa/opencl/actor_facade.hpp"
#include "cppa/detail/singleton_mixin.hpp"
#include "cppa/detail/singleton_manager.hpp"
#include "cppa/intrusive/blocking_single_reader_queue.hpp"
namespace cppa { namespace opencl {
#ifdef CPPA_OPENCL
class command_dispatcher {
struct worker;
friend struct worker;
friend class detail::singleton_manager;
friend class program;
friend void enqueue_to_dispatcher(command_dispatcher*, command*);
public:
void enqueue();
template<typename Ret, typename... Args>
actor_ptr spawn(const program& prog,
const std::string& kernel_name) {
return new actor_facade<Ret (Args...)>(this, prog, kernel_name);
}
template<typename Ret, typename... Args>
actor_ptr spawn(const std::string& kernel_source,
const std::string& kernel_name) {
return spawn<Ret, Args...>(program{kernel_source}, kernel_name);
}
private:
struct device_info {
unsigned id;
command_queue_ptr cmd_queue;
device_ptr dev_id;
size_t max_itms_per_grp;
cl_uint max_dim;
std::vector<size_t> 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,
std::vector<size_t> 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(std::move(max_itms_per_dim)) { }
};
typedef intrusive::blocking_single_reader_queue<command> job_queue;
static inline command_dispatcher* create_singleton() {
return new command_dispatcher;
}
inline void initialize() { }
inline void dispose() { delete this; }
inline void destroy() { delete this; }
void initialize();
void dispose();
void destroy();
std::atomic<unsigned> dev_id_gen;
job_queue m_job_queue;
std::thread m_supervisor;
std::vector<device_info> m_devices;
context_ptr m_context;
static void worker_loop(worker*);
static void supervisor_loop(command_dispatcher *scheduler, job_queue*);
};
#else // CPPA_OPENCL
class command_dispatcher : public detail::singleton_mixin<command_dispatcher> {
......
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \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 CPPA_OPENCL_GLOBAL_HPP
#define CPPA_OPENCL_GLOBAL_HPP
#include <string>
#if defined __APPLE__ || defined(MACOSX)
#include <OpenCL/opencl.h>
#else
#include <CL/opencl.h>
#endif
namespace cppa { namespace opencl {
std::string get_opencl_error(cl_int err);
} } // namespace cppa::opencl
#endif // CPPA_OPENCL_GLOBAL_HPP
......@@ -32,14 +32,32 @@
#ifndef CPPA_OPENCL_PROGRAM_HPP
#define CPPA_OPENCL_PROGRAM_HPP
#include <memory>
#include "cppa/opencl/global.hpp"
#include "cppa/opencl/smart_ptr.hpp"
namespace cppa { namespace opencl {
template<typename Signature>
class actor_facade;
class program {
template<typename Signature>
friend class actor_facade;
public:
program(const std::string& kernel_source);
private:
context_ptr m_context;
program_ptr m_program;
};
} // namespace cppa::opencl
} } // namespace cppa::opencl
#endif // CPPA_OPENCL_PROGRAM_HPP
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \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 SMART_PTR_HPP
#define SMART_PTR_HPP
#include <memory>
#include <algorithm>
#include <type_traits>
namespace cppa { namespace opencl {
template<typename T, cl_int (*ref)(T), cl_int (*deref)(T)>
class smart_ptr {
typedef typename std::remove_pointer<T>::type element_type;
typedef element_type* pointer;
typedef element_type& reference;
typedef const element_type* const_pointer;
typedef const element_type& const_reference;
public:
smart_ptr(pointer ptr = nullptr) : m_ptr(ptr) {
if (m_ptr) ref(m_ptr);
}
~smart_ptr() { reset(); }
smart_ptr(const smart_ptr& other) : m_ptr(other.m_ptr) {
if (m_ptr) ref(m_ptr);
}
smart_ptr(smart_ptr&& other) : m_ptr(other.m_ptr) {
other.m_ptr = nullptr;
}
smart_ptr& operator=(pointer ptr) {
reset(ptr);
return *this;
}
smart_ptr& operator=(smart_ptr&& other) {
std::swap(m_ptr, other.m_ptr);
return *this;
}
smart_ptr& operator=(const smart_ptr& other) {
smart_ptr tmp{other};
std::swap(m_ptr, tmp.m_ptr);
return *this;
}
inline void reset(pointer ptr = nullptr) {
if (m_ptr) deref(m_ptr);
m_ptr = ptr;
if (ptr) ref(ptr);
}
// does not modify reference count of ptr
inline void adopt(pointer ptr) {
reset();
m_ptr = ptr;
}
inline pointer get() const { return m_ptr; }
inline pointer operator->() const { return m_ptr; }
inline reference operator*() const { return *m_ptr; }
inline bool operator!() const { return m_ptr == nullptr; }
inline explicit operator bool() const { return m_ptr != nullptr; }
private:
pointer m_ptr;
};
typedef smart_ptr<cl_mem,clRetainMemObject,clReleaseMemObject> mem_ptr;
typedef smart_ptr<cl_kernel,clRetainKernel,clReleaseKernel> kernel_ptr;
typedef smart_ptr<cl_device_id,clRetainDevice,clReleaseDevice> device_ptr;
typedef smart_ptr<cl_context,clRetainContext,clReleaseContext> context_ptr;
typedef smart_ptr<cl_program,clRetainProgram,clReleaseProgram> program_ptr;
typedef smart_ptr<cl_command_queue,clRetainCommandQueue,clReleaseCommandQueue>
command_queue_ptr;
} } // namespace cppa::opencl
#endif // SMART_PTR_HPP
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \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/>. *
\******************************************************************************/
#include "cppa/opencl/actor_facade.hpp"
#include "cppa/opencl/command_dispatcher.hpp"
namespace cppa { namespace opencl {
void enqueue_to_dispatcher(command_dispatcher* dispatcher,
command* cmd) {
dispatcher->m_job_queue.push_back(cmd);
}
} } // namespace cppa::opencl
......@@ -28,5 +28,206 @@
* along with libcppa. If not, see <http://www.gnu.org/licenses/>. *
\******************************************************************************/
#include <sstream>
#include <iostream>
#include <stdexcept>
#include "cppa/cppa.hpp"
#include "cppa/opencl/command_dispatcher.hpp"
using namespace std;
namespace cppa { namespace opencl {
struct command_dispatcher::worker {
command_dispatcher* m_parent;
typedef unique_ptr<command> job_ptr;
job_queue* m_job_queue;
thread m_thread;
worker(command_dispatcher* parent, job_queue* jq)
: m_parent(parent), m_job_queue(jq) { }
void start() {
m_thread = thread(&command_dispatcher::worker_loop, this);
}
worker(const worker&) = delete;
worker& operator=(const worker&) = delete;
void operator()() {
job_ptr job;
for (;;) {
/* wait for device */
/* get results */
/* wait for job */
job.reset(m_job_queue->pop());
if(job->fun) {
try {
cl_command_queue cmd_q =
m_parent->m_devices.front().cmd_queue.get();
job->fun(cmd_q);
}
catch (exception& e) {
cerr << e.what() << endl;
}
}
else {
cout << "worker done" << endl;
return;
}
// ...?
}
}
};
void command_dispatcher::worker_loop(command_dispatcher::worker* w) {
(*w)();
}
void command_dispatcher::supervisor_loop(command_dispatcher* scheduler,
job_queue* jq) {
unique_ptr<command_dispatcher::worker> worker;
worker.reset(new command_dispatcher::worker(scheduler, jq));
worker->start();
worker->m_thread.join();
worker.reset();
cout << "supervisor done" << endl;
}
void command_dispatcher::initialize() {
cl_int err{0};
/* find up to two available platforms */
vector<cl_platform_id> ids(2);
cl_uint number_of_platforms;
err = clGetPlatformIDs(ids.size(), ids.data(), &number_of_platforms);
if (err != CL_SUCCESS) {
throw logic_error("[!!!] clGetPlatformIDs: '"
+ get_opencl_error(err)
+ "'.");
}
else if (number_of_platforms < 1) {
throw logic_error("[!!!] clGetPlatformIDs: 'no platforms found'.");
}
/* find gpu devices on our platform */
int pid{0};
cl_uint num_devices{0};
cout << "Currently only looking for cpu devices!" << endl;
cl_device_type dev_type{CL_DEVICE_TYPE_CPU /*CL_DEVICE_TYPE_GPU*/};
err = clGetDeviceIDs(ids[pid], dev_type, 0, NULL, &num_devices);
if (err == CL_DEVICE_NOT_FOUND) {
cout << "NO GPU DEVICES FOUND! LOOKING FOR CPU DEVICES NOW ..." << endl;
dev_type = CL_DEVICE_TYPE_CPU;
err = clGetDeviceIDs(ids[pid], dev_type, 0, NULL, &num_devices);
}
if (err != CL_SUCCESS) {
throw runtime_error("[!!!] clGetDeviceIDs: '"
+ get_opencl_error(err)
+ "'.");
}
vector<cl_device_id> devices(num_devices);
err = clGetDeviceIDs(ids[pid], dev_type, num_devices, devices.data(), NULL);
if (err != CL_SUCCESS) {
throw runtime_error("[!!!] clGetDeviceIDs: '"
+ get_opencl_error(err)
+ "'.");
}
/* create a context */
m_context.adopt(clCreateContext(0, 1, devices.data(), NULL, NULL, &err));
if (err != CL_SUCCESS) {
throw runtime_error("[!!!] clCreateContext: '"
+ get_opencl_error(err)
+ "'.");
}
for (auto& d : devices) {
device_ptr device;
device.adopt(d);
unsigned id{++dev_id_gen};
command_queue_ptr cmd_queue;
cmd_queue.adopt(clCreateCommandQueue(m_context.get(),
device.get(),
CL_QUEUE_PROFILING_ENABLE,
&err));
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "[!!!] clCreateCommandQueue (" << id << "): '"
<< get_opencl_error(err) << "'.";
throw runtime_error(oss.str());
}
size_t return_size{0};
size_t max_work_group_size{0};
err = clGetDeviceInfo(device.get(),
CL_DEVICE_MAX_WORK_GROUP_SIZE,
sizeof(size_t),
&max_work_group_size,
&return_size);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "[!!!] clGetDeviceInfo ("
<< id
<< ":CL_DEVICE_MAX_WORK_GROUP_SIZE): '"
<< get_opencl_error(err) << "'.";
throw runtime_error(oss.str());
}
cl_uint max_work_item_dimensions = 0;
err = clGetDeviceInfo(device.get(),
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS,
sizeof(cl_uint),
&max_work_item_dimensions,
&return_size);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "[!!!] clGetDeviceInfo ("
<< id
<< ":CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): '"
<< get_opencl_error(err) << "'.";
throw runtime_error(oss.str());
}
vector<size_t> max_work_items_per_dim(max_work_item_dimensions);
err = clGetDeviceInfo(device.get(),
CL_DEVICE_MAX_WORK_ITEM_SIZES,
sizeof(size_t)*max_work_item_dimensions,
max_work_items_per_dim.data(),
&return_size);
if (err != CL_SUCCESS) {
ostringstream oss;
oss << "[!!!] clGetDeviceInfo ("
<< id
<< ":CL_DEVICE_MAX_WORK_ITEM_SIZES): '"
<< get_opencl_error(err) << "'.";
throw runtime_error(oss.str());
}
device_info dev_info{id,
cmd_queue,
device,
max_work_group_size,
max_work_item_dimensions,
move(max_work_items_per_dim)};
m_devices.push_back(move(dev_info));
}
m_supervisor = thread(&command_dispatcher::supervisor_loop,
this,
&m_job_queue);
}
void command_dispatcher::destroy() {
m_job_queue.push_back(new command);
m_supervisor.join();
delete this;
}
void command_dispatcher::dispose() {
delete this;
}
} } // namespace cppa::opencl
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \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/>. *
\******************************************************************************/
#include "cppa/opencl/global.hpp"
namespace cppa { namespace opencl {
std::string get_opencl_error(cl_int err) {
switch (err) {
case CL_SUCCESS:
return "CL_SUCCESS";
case CL_DEVICE_NOT_FOUND:
return "CL_DEVICE_NOT_FOUND";
case CL_DEVICE_NOT_AVAILABLE:
return "CL_DEVICE_NOT_AVAILABLE";
case CL_COMPILER_NOT_AVAILABLE:
return "CL_COMPILER_NOT_AVAILABLE";
case CL_MEM_OBJECT_ALLOCATION_FAILURE:
return "CL_MEM_OBJECT_ALLOCATION_FAILURE";
case CL_OUT_OF_RESOURCES:
return "CL_OUT_OF_RESOURCES";
case CL_OUT_OF_HOST_MEMORY:
return "CL_OUT_OF_HOST_MEMORY";
case CL_PROFILING_INFO_NOT_AVAILABLE:
return "CL_PROFILING_INFO_NOT_AVAILABLE";
case CL_MEM_COPY_OVERLAP:
return "CL_MEM_COPY_OVERLAP";
case CL_IMAGE_FORMAT_MISMATCH:
return "CL_IMAGE_FORMAT_MISMATCH";
case CL_IMAGE_FORMAT_NOT_SUPPORTED:
return "CL_IMAGE_FORMAT_NOT_SUPPORTED";
case CL_BUILD_PROGRAM_FAILURE:
return "CL_BUILD_PROGRAM_FAILURE";
case CL_MAP_FAILURE:
return "CL_MAP_FAILURE";
case CL_INVALID_VALUE:
return "CL_INVALID_VALUE";
case CL_INVALID_DEVICE_TYPE:
return "CL_INVALID_DEVICE_TYPE";
case CL_INVALID_PLATFORM:
return "CL_INVALID_PLATFORM";
case CL_INVALID_DEVICE:
return "CL_INVALID_DEVICE";
case CL_INVALID_CONTEXT:
return "CL_INVALID_CONTEXT";
case CL_INVALID_QUEUE_PROPERTIES:
return "CL_INVALID_QUEUE_PROPERTIES";
case CL_INVALID_COMMAND_QUEUE:
return "CL_INVALID_COMMAND_QUEUE";
case CL_INVALID_HOST_PTR:
return "CL_INVALID_HOST_PTR";
case CL_INVALID_MEM_OBJECT:
return "CL_INVALID_MEM_OBJECT";
case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR:
return "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR";
case CL_INVALID_IMAGE_SIZE:
return "CL_INVALID_IMAGE_SIZE";
case CL_INVALID_SAMPLER:
return "CL_INVALID_SAMPLER";
case CL_INVALID_BINARY:
return "CL_INVALID_BINARY";
case CL_INVALID_BUILD_OPTIONS:
return "CL_INVALID_BUILD_OPTIONS";
case CL_INVALID_PROGRAM:
return "CL_INVALID_PROGRAM";
case CL_INVALID_PROGRAM_EXECUTABLE:
return "CL_INVALID_PROGRAM_EXECUTABLE";
case CL_INVALID_KERNEL_NAME:
return "CL_INVALID_KERNEL_NAME";
case CL_INVALID_KERNEL_DEFINITION:
return "CL_INVALID_KERNEL_DEFINITION";
case CL_INVALID_KERNEL:
return "CL_INVALID_KERNEL";
case CL_INVALID_ARG_INDEX:
return "CL_INVALID_ARG_INDEX";
case CL_INVALID_ARG_VALUE:
return "CL_INVALID_ARG_VALUE";
case CL_INVALID_ARG_SIZE:
return "CL_INVALID_ARG_SIZE";
case CL_INVALID_KERNEL_ARGS:
return "CL_INVALID_KERNEL_ARGS";
case CL_INVALID_WORK_DIMENSION:
return "CL_INVALID_WORK_DIMENSION";
case CL_INVALID_WORK_GROUP_SIZE:
return "CL_INVALID_WORK_GROUP_SIZE";
case CL_INVALID_WORK_ITEM_SIZE:
return "CL_INVALID_WORK_ITEM_SIZE";
case CL_INVALID_GLOBAL_OFFSET:
return "CL_INVALID_GLOBAL_OFFSET";
case CL_INVALID_EVENT_WAIT_LIST:
return "CL_INVALID_EVENT_WAIT_LIST";
case CL_INVALID_EVENT:
return "CL_INVALID_EVENT";
case CL_INVALID_OPERATION:
return "CL_INVALID_OPERATION";
case CL_INVALID_GL_OBJECT:
return "CL_INVALID_GL_OBJECT";
case CL_INVALID_BUFFER_SIZE:
return "CL_INVALID_BUFFER_SIZE";
case CL_INVALID_MIP_LEVEL:
return "CL_INVALID_MIP_LEVEL";
case CL_INVALID_GLOBAL_WORK_SIZE:
return "CL_INVALID_GLOBAL_WORK_SIZE";
default: return "UNKNOWN_ERROR";
}
}
} } // namespace cppa::opencl
/******************************************************************************\
* ___ __ *
* /\_ \ __/\ \ *
* \//\ \ /\_\ \ \____ ___ _____ _____ __ *
* \ \ \ \/\ \ \ '__`\ /'___\/\ '__`\/\ '__`\ /'__`\ *
* \_\ \_\ \ \ \ \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/>. *
\******************************************************************************/
#include "cppa/opencl/program.hpp"
#include "cppa/opencl/command_dispatcher.hpp"
namespace cppa { namespace opencl {
program::program(const std::string& kernel_source) {
m_context =
cppa::detail::singleton_manager::get_command_dispatcher()->m_context;
cl_int err{0};
/* create program object from kernel source */
size_t kernel_source_length = kernel_source.size();
const char *kernel_source_cstr = kernel_source.c_str();
m_program.adopt(clCreateProgramWithSource(m_context.get(),
1,
&kernel_source_cstr,
&kernel_source_length,
&err));
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clCreateProgramWithSource: '"
+ get_opencl_error(err)
+ "'.");
}
/* build programm from program object */
err = clBuildProgram(m_program.get(), 0, NULL, NULL, NULL, NULL);
if (err != CL_SUCCESS) {
throw std::runtime_error("[!!!] clBuildProgram: '"
+ get_opencl_error(err)
+ "'.");
}
}
} } // 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