Commit 318bf45c authored by Dominik Charousset's avatar Dominik Charousset

Merge branch 'topic/opencl' into unstable

parents 7f011f0b d76bbd6c
......@@ -303,3 +303,5 @@ unit_testing/test_sync_send.cpp
unit_testing/test_tuple.cpp
unit_testing/test_uniform_type.cpp
unit_testing/test_yield_interface.cpp
examples/opencl/simple_matrix.cpp
examples/opencl/proper_matrix.cpp
......@@ -56,7 +56,7 @@ template<typename MapArgs, typename MapResult>
struct get_cl_spawn_helper;
template<typename R, typename... Ts>
struct get_cl_spawn_helper<std::function<option<cow_tuple<Ts...>> (const any_tuple&)>,
struct get_cl_spawn_helper<std::function<option<cow_tuple<Ts...>> (any_tuple)>,
std::function<any_tuple (R&)>> {
typedef cl_spawn_helper<R (const Ts&...)> type;
};
......@@ -91,6 +91,36 @@ actor_ptr spawn_cl(const char* source,
f0{map_args}, f1{map_result});
}
template<typename Signature, typename... Ts>
actor_ptr spawn_cl(const opencl::program& prog,
const char* fun_name,
std::vector<size_t> dimensions,
std::vector<size_t> offset = {},
std::vector<size_t> local_dims = {}) {
using std::move;
cl_spawn_helper<Signature> f;
return f(prog, fun_name, move(dimensions), move(offset), move(local_dims));
}
template<typename MapArgs, typename MapResult>
actor_ptr spawn_cl(const opencl::program& prog,
const char* fun_name,
MapArgs map_args,
MapResult map_result,
std::vector<size_t> dimensions,
std::vector<size_t> offset = {},
std::vector<size_t> local_dims = {}) {
using std::move;
typedef typename util::get_callable_trait<MapArgs>::type t0;
typedef typename util::get_callable_trait<MapResult>::type t1;
typedef typename t0::fun_type f0;
typedef typename t1::fun_type f1;
typename get_cl_spawn_helper<f0,f1>::type f;
return f(prog, fun_name,
move(dimensions), move(offset), move(local_dims),
f0{map_args}, f1{map_result});
}
} // namespace cppa
#endif // CPPA_OPENCL_HPP
......@@ -69,7 +69,7 @@ class actor_facade<Ret(Args...)> : public actor {
public:
typedef cow_tuple<typename util::rm_ref<Args>::type...> args_tuple;
typedef std::function<option<args_tuple>(const any_tuple&)> arg_mapping;
typedef std::function<option<args_tuple>(any_tuple)> arg_mapping;
typedef std::function<any_tuple(Ret&)> result_mapping;
static actor_facade* create(command_dispatcher* dispatcher,
......@@ -152,7 +152,7 @@ class actor_facade<Ret(Args...)> : public actor {
void enqueue_impl(const actor_ptr& sender, any_tuple msg, message_id id, util::int_list<Is...>) {
auto opt = m_map_args(msg);
if (opt) {
response_handle handle{this, sender, id};
response_handle handle{this, sender, id.response_id()};
size_t number_of_values{1};
std::for_each(m_global_dimensions.begin(),
m_global_dimensions.end(),
......
......@@ -148,13 +148,13 @@ class command_impl : public command {
/* get results from gpu */
cl_int err{0};
cl_event read_event;
T results(m_number_of_values);
T result(m_number_of_values);
err = clEnqueueReadBuffer(m_queue.get(),
m_arguments[0].get(),
CL_TRUE,
0,
sizeof(typename T::value_type) * m_number_of_values,
results.data(),
result.data(),
0,
NULL,
&read_event);
......@@ -164,7 +164,7 @@ class command_impl : public command {
+ get_opencl_error(err)
+ "'.");
}
auto mapped_result = m_map_result(results);
auto mapped_result = m_map_result(result);
reply_tuple_to(m_handle, mapped_result);
//reply_to(m_handle, results);
}
......
......@@ -87,7 +87,7 @@ class command_dispatcher {
std::vector<size_t> global_dims,
std::vector<size_t> global_offs,
std::vector<size_t> local_dims,
std::function<option<cow_tuple<typename util::rm_ref<Args>::type...>>(const any_tuple&)> map_args,
std::function<option<cow_tuple<typename util::rm_ref<Args>::type...>>(any_tuple)> map_args,
std::function<any_tuple(Ret&)> map_result)
{
return actor_facade<Ret (Args...)>::create(this,
......
......@@ -46,3 +46,14 @@ if (NOT "${CPPA_NO_QT_EXAMPLES}" STREQUAL "yes")
add_dependencies(qt_group_chat all_examples)
endif (QT4_FOUND)
endif()
if (ENABLE_OPENCL)
macro(add_opencl name folder)
add_executable(${name} ${folder}/${name}.cpp ${ARGN})
target_link_libraries(${name} ${CMAKE_DL_LIBS} ${CPPA_LIBRARY} ${PTHREAD_LIBRARIES} ${OPENCL_LIBRARIES})
add_dependencies(${name} all_examples)
endmacro()
add_opencl(simple_matrix opencl)
add_opencl(proper_matrix opencl)
endif (ENABLE_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 <vector>
#include <iomanip>
#include <numeric>
#include <iostream>
#include "cppa/cppa.hpp"
#include "cppa/opencl.hpp"
using namespace std;
using namespace cppa;
static constexpr size_t matrix_size = 8;
static constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2
namespace { constexpr const char* kernel_source = R"__(
__kernel void matrix_mult(__global float* matrix1,
__global float* matrix2,
__global float* output) {
int size = get_global_size(0); // == get_global_size_(1);
int x = get_global_id(0); // get position in matrix
int y = get_global_id(1); // this work item should calculate
int idx = 0;
float result = 0;
while (idx < size) { // calculate one element of the matrix
float i = matrix1[idx+y*size];
float j = matrix2[x+idx*size];
float tmp = i*j;
result = result + tmp;
++idx;
}
output[x+y*size] = result;
}
)__"; }
// represensts a square matrix
template<size_t Size>
class square_matrix {
public:
static constexpr size_t num_elements = Size * Size;
square_matrix(square_matrix&&) = default;
square_matrix(const square_matrix&) = default;
square_matrix& operator=(square_matrix&&) = default;
square_matrix& operator=(const square_matrix&) = default;
square_matrix() {
m_data.resize(num_elements);
}
square_matrix(vector<float> d) : m_data(std::move(d)) { }
square_matrix(const std::initializer_list<float>& args) : m_data(args) {
m_data.resize(num_elements);
}
inline float& operator()(size_t row, size_t column) {
return m_data[row + column * Size];
}
inline const float& operator()(size_t row, size_t column) const {
return m_data[row + column * Size];
}
inline void zeroize() {
std::fill(m_data.begin(), m_data.end(), 0);
}
inline void iota_fill() {
std::iota(m_data.begin(), m_data.end(), 0);
}
typedef typename vector<float>::const_iterator const_iterator;
const_iterator begin() const { return m_data.begin(); }
const_iterator end() const { return m_data.end(); }
vector<float>& data() { return m_data; }
const vector<float>& data() const { return m_data; }
private:
vector<float> m_data;
};
// to annouce the square_matrix to lobccpa
// these operaters have to be implemented
template<size_t Size>
inline bool operator==(const square_matrix<Size>& lhs, const square_matrix<Size>& rhs) {
return std::equal(lhs.begin(), lhs.end(), rhs.begin());
}
template<size_t Size>
inline bool operator!=(const square_matrix<Size>& lhs, const square_matrix<Size>& rhs) {
return !(lhs == rhs);
}
using matrix_type = square_matrix<matrix_size>;
void multiplier() {
// create two matrix and fill them with values 0, 1, 2, 3, 4, ...
matrix_type m1;
matrix_type m2;
m1.iota_fill();
m2.iota_fill();
// print "source" matrix
cout << "calculating square of matrix:" << endl;
for (size_t y = 0; y < matrix_size; ++y) {
for (size_t x = 0; x < matrix_size; ++x) {
cout << fixed << setprecision(2) << setw(6) << m1(x,y);
}
cout << endl;
}
cout << endl;
// spawn an opencl actor
// 1st arg: source code of one or more opencl kernels
// 2nd arg: name of the kernel to use
auto worker = spawn_cl(kernel_source, kernel_name,
// 3rd arg: the opencl actor can only handle vectors,
// this function creates vectors from a tuple of matrix_types
[] (any_tuple msg) -> option<cow_tuple<vector<float>,vector<float>>> {
auto opt = tuple_cast<matrix_type,matrix_type>(msg);
if (opt) {
return {move(get_ref<0>(*opt).data()),
move(get_ref<1>(*opt).data())};
}
return {};
},
// 4th arg: builds a matrix_type from a vector
// allows the actor to pass a matrix_type back instead of a vector
[] (vector<float>& result) -> any_tuple {
return make_any_tuple<matrix_type>(move(result));
},
// 5th arg: global dimension arguments for opencl's enqueue
// creates matrix_size * matrix_size global work items
{matrix_size, matrix_size}
);
// send both matrices to the actor and
// wait for results in form of a matrix_type
sync_send(worker, move(m1), move(m2)).then(
[](const matrix_type& result) {
cout << "result:" << endl;
for (size_t y = 0; y < matrix_size; ++y) {
for (size_t x = 0; x < matrix_size; ++x) {
cout << fixed << setprecision(2) << setw(9) << result(x,y);
}
cout << endl;
}
}
);
}
int main() {
// matrix_type ist not a simple type,
// it must be annouced to libcppa
announce<matrix_type>();
spawn(multiplier);
await_all_others_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 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 <vector>
#include <iomanip>
#include <numeric>
#include <iostream>
#include "cppa/cppa.hpp"
#include "cppa/opencl.hpp"
using namespace std;
using namespace cppa;
static constexpr size_t matrix_size = 8;
static constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2
namespace { constexpr const char* kernel_source = R"__(
__kernel void matrix_mult(__global float* matrix1,
__global float* matrix2,
__global float* output) {
int size = get_global_size(0); // == get_global_size_(1);
int x = get_global_id(0);
int y = get_global_id(1);
int idx = 0;
float result = 0;
while (idx < size) {
float i = matrix1[idx+y*size];
float j = matrix2[x+idx*size];
float tmp = i*j;
result = result + tmp;
++idx;
}
output[x+y*size] = result;
}
)__"; }
void multiplier() {
// the opencl actor only understands vectors
// so these vectors represent the matrices
vector<float> m1(matrix_size * matrix_size);
vector<float> m2(matrix_size * matrix_size);
// fill each with values from 0 to matix_size * matrix_size - 1
iota(m1.begin(), m1.end(), 0);
iota(m2.begin(), m2.end(), 0);
// print "source" matrix
cout << "calculating square of matrix:" << endl;
for (size_t y = 0; y < matrix_size; ++y) {
for (size_t x = 0; x < matrix_size; ++x) {
cout << fixed << setprecision(2) << setw(6) << m1[x+y*matrix_size];
}
cout << endl;
}
cout << endl;
// spawn an opencl actor
// 1st arg: sourec code of one or more kernels
// 2nd arg: name of the kernel to use
// 3rd arg: global dimension arguments for opencl's enqueue
// creates matrix_size * matrix_size global work items
auto worker =
spawn_cl<vector<float>(vector<float>&,vector<float>&)>(kernel_source,
kernel_name,
{matrix_size, matrix_size});
// send both matrices to the actor and wait for a result
sync_send(worker, move(m1), move(m2)).then(
[](const vector<float>& result) {
cout << "result:" << endl;
for (size_t column = 0; column < matrix_size; ++column) {
for (size_t row = 0; row < matrix_size; ++row) {
cout << fixed << setprecision(2) << setw(9)
<< result[row+column*matrix_size];
}
cout << endl;
}
}
);
}
int main() {
announce<vector<float>>();
spawn(multiplier);
await_all_others_done();
shutdown();
return 0;
}
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