Commit 4c96b7be authored by Marian Triebe's avatar Marian Triebe

Add opencl examples, fix linker problem

parent d1abaf6a
...@@ -5,36 +5,35 @@ project(caf_opencl C CXX) ...@@ -5,36 +5,35 @@ project(caf_opencl C CXX)
# e.g., for creating proper Xcode projects # e.g., for creating proper Xcode projects
file(GLOB LIBCAF_OPENCL_HDRS "caf/opencl/*.hpp") file(GLOB LIBCAF_OPENCL_HDRS "caf/opencl/*.hpp")
# list cpp files excluding platform-dependent files if (OPENCL_LIBRARIES)
set (LIBCAF_OPENCL_SRCS # list cpp files excluding platform-dependent files
src/global.cpp set (LIBCAF_OPENCL_SRCS
src/opencl_metainfo.cpp src/global.cpp
src/program.cpp) src/opencl_metainfo.cpp
src/program.cpp)
# build shared library if not compiling static only # build shared library if not compiling static only
if (NOT "${CAF_BUILD_STATIC_ONLY}" STREQUAL "yes") if(NOT "${CAF_BUILD_STATIC_ONLY}" STREQUAL "yes")
add_library(libcaf_opencl SHARED ${LIBCAF_OPENCL_SRCS} ${LIBCAF_OPENCL_HDRS}) add_library(libcaf_opencl SHARED ${LIBCAF_OPENCL_SRCS} ${LIBCAF_OPENCL_HDRS})
target_link_libraries(libcaf_opencl ${LD_FLAGS} ${LIBCAF_CORE_LIBRARY}) target_link_libraries(libcaf_opencl ${LD_FLAGS} ${LIBCAF_CORE_LIBRARY})
set(LIBRARY_SOVERSION ${CAF_VERSION_MAJOR}) set(LIBRARY_SOVERSION ${CAF_VERSION_MAJOR})
set_target_properties(libcaf_opencl set_target_properties(libcaf_opencl
PROPERTIES PROPERTIES
SOVERSION ${LIBRARY_SOVERSION} SOVERSION ${LIBRARY_SOVERSION}
VERSION ${CAF_VERSION} VERSION ${CAF_VERSION}
OUTPUT_NAME caf_opencl) OUTPUT_NAME caf_opencl)
if (NOT MINGW) if(NOT MINGW)
install(TARGETS libcaf_opencl LIBRARY DESTINATION lib) install(TARGETS libcaf_opencl LIBRARY DESTINATION lib)
endif()
endif() endif()
endif () # build static library only if --build-static or --build-static-only was set
if("${CAF_BUILD_STATIC_ONLY}" STREQUAL "yes" OR "${CAF_BUILD_STATIC}" STREQUAL "yes")
# build static library only if --build-static or --build-static-only was set add_library(libcaf_openclStatic STATIC ${LIBCAF_OPENCL_HDRS} ${LIBCAF_OPENCL_SRCS})
if ("${CAF_BUILD_STATIC_ONLY}" STREQUAL "yes" OR "${CAF_BUILD_STATIC}" STREQUAL "yes") set_target_properties(libcaf_openclStatic PROPERTIES OUTPUT_NAME caf_opencl_static)
add_library(libcaf_openclStatic STATIC ${LIBCAF_OPENCL_HDRS} ${LIBCAF_OPENCL_SRCS}) install(TARGETS libcaf_openclStatic ARCHIVE DESTINATION lib)
set_target_properties(libcaf_openclStatic PROPERTIES OUTPUT_NAME caf_opencl_static) endif()
install(TARGETS libcaf_openclStatic ARCHIVE DESTINATION lib) link_directories(${LD_DIRS})
endif () include_directories(. ${INCLUDE_DIRS})
endif()
link_directories(${LD_DIRS})
include_directories(. ${INCLUDE_DIRS})
# install includes # install includes
install(DIRECTORY caf/ DESTINATION include/caf FILES_MATCHING PATTERN "*.hpp") install(DIRECTORY caf/ DESTINATION include/caf FILES_MATCHING PATTERN "*.hpp")
......
...@@ -101,7 +101,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor { ...@@ -101,7 +101,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
} }
void enqueue(const actor_addr &sender, message_id mid, message content, void enqueue(const actor_addr &sender, message_id mid, message content,
execution_unit *host) override { execution_unit*) override {
CAF_LOG_TRACE(""); CAF_LOG_TRACE("");
typename detail::il_indices<detail::type_list<Args...>>::type indices; typename detail::il_indices<detail::type_list<Args...>>::type indices;
enqueue_impl(sender, mid, std::move(content), indices); enqueue_impl(sender, mid, std::move(content), indices);
......
...@@ -55,9 +55,9 @@ class opencl_metainfo : public detail::abstract_singleton { ...@@ -55,9 +55,9 @@ class opencl_metainfo : public detail::abstract_singleton {
static opencl_metainfo* instance(); static opencl_metainfo* instance();
private: private:
opencl_metainfo(); opencl_metainfo() = default;
void stop(); void stop() override;
void initialize() override; void initialize() override;
void dispose() override; void dispose() override;
......
...@@ -23,10 +23,11 @@ ...@@ -23,10 +23,11 @@
#include <algorithm> #include <algorithm>
#include <functional> #include <functional>
#include "caf/actor_cast.hpp"
#include "caf/optional.hpp" #include "caf/optional.hpp"
#include "cppa/cow_tuple.hpp" #include "cppa/cow_tuple.hpp"
#include "caf/actor_cast.hpp" #include "caf/detail/type_traits.hpp"
#include "caf/detail/limited_vector.hpp" #include "caf/detail/limited_vector.hpp"
#include "caf/opencl/global.hpp" #include "caf/opencl/global.hpp"
...@@ -76,7 +77,7 @@ struct cl_spawn_helper<R(Ts...), void> { ...@@ -76,7 +77,7 @@ struct cl_spawn_helper<R(Ts...), void> {
using std::move; using std::move;
using std::forward; using std::forward;
map_arg_fun f0 = [](message msg) { map_arg_fun f0 = [](message msg) {
return tuple_cast<typename util::rm_const_and_ref< return tuple_cast<typename detail::rm_const_and_ref<
typename carr_to_vec<Ts>::type>::type...>(msg); typename carr_to_vec<Ts>::type>::type...>(msg);
}; };
map_res_fun f1 = [](result_type& result) { map_res_fun f1 = [](result_type& result) {
...@@ -144,8 +145,8 @@ spawn_cl(const opencl::program& prog, const char* fname, MapArgs map_args, ...@@ -144,8 +145,8 @@ spawn_cl(const opencl::program& prog, const char* fname, MapArgs map_args,
const opencl::dim_vec& offset = {}, const opencl::dim_vec& offset = {},
const opencl::dim_vec& local_dims = {}, size_t result_size = 0) { const opencl::dim_vec& local_dims = {}, size_t result_size = 0) {
using std::move; using std::move;
using f0 = typename util::get_callable_trait<MapArgs>::fun_type; using f0 = typename detail::get_callable_trait<MapArgs>::fun_type;
using f1 = typename util::get_callable_trait<MapResult>::fun_type; using f1 = typename detail::get_callable_trait<MapResult>::fun_type;
detail::cl_spawn_helper<f0, f1> f; detail::cl_spawn_helper<f0, f1> f;
return f(f0{move(map_args)}, f1{move(map_result)}, prog, fname, dims, offset, return f(f0{move(map_args)}, f1{move(map_result)}, prog, fname, dims, offset,
local_dims, result_size); local_dims, result_size);
......
cmake_minimum_required(VERSION 2.8)
project(caf_examples_opencl CXX)
if(OPENCL_LIBRARIES)
add_custom_target(opencl_examples)
include_directories(${LIBCAF_INCLUDE_DIRS})
if(${CMAKE_SYSTEM_NAME} MATCHES "Window")
set(WSLIB -lws2_32)
else()
set(WSLIB)
endif()
macro(add name folder)
add_executable(${name} ${folder}/${name}.cpp ${ARGN})
target_link_libraries(${name}
${LD_FLAGS}
${LIBCAF_LIBRARIES}
${PTHREAD_LIBRARIES}
${WSLIB}
${OPENCL_LIBRARIES})
add_dependencies(${name} all_examples)
endmacro()
add(proper_matrix .)
add(simple_matrix .)
endif()
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2014 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENCE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include <vector>
#include <iomanip>
#include <numeric>
#include <cassert>
#include <iostream>
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
using namespace std;
using namespace caf;
namespace {
using fvec = vector<float>;
constexpr size_t matrix_size = 8;
constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2
// last parameter is, by convention, the output parameter
constexpr const char* kernel_source = R"__(
__kernel void matrix_mult(__global float* matrix1,
__global float* matrix2,
__global float* output) {
// we only use square matrices, hence: width == height
size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0);
size_t y = get_global_id(1);
float result = 0;
for (size_t idx = 0; idx < size; ++idx) {
result += matrix1[idx + y * size] * matrix2[x + idx * size];
}
output[x+y*size] = result;
}
)__";
} // namespace <anonymous>
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(num_elements) { }
explicit square_matrix(fvec d) : m_data(move(d)) {
assert(m_data.size() == num_elements);
}
inline float& operator()(size_t column, size_t row) {
return m_data[column + row * Size];
}
inline const float& operator()(size_t column, size_t row) const {
return m_data[column + row * Size];
}
inline void iota_fill() {
iota(m_data.begin(), m_data.end(), 0);
}
typedef typename fvec::const_iterator const_iterator;
const_iterator begin() const { return m_data.begin(); }
const_iterator end() const { return m_data.end(); }
fvec& data() { return m_data; }
const fvec& data() const { return m_data; }
private:
fvec m_data;
};
template<size_t Size>
string to_string(const square_matrix<Size>& m) {
ostringstream oss;
oss.fill(' ');
for (size_t row = 0; row < Size; ++row) {
for (size_t column = 0; column < Size; ++column) {
oss << fixed << setprecision(2) << setw(9) << m(column, row);
}
oss << '\n';
}
return oss.str();
}
// to annouce the square_matrix to libcaf, we have to define these operators
template<size_t Size>
inline bool operator==(const square_matrix<Size>& lhs,
const square_matrix<Size>& rhs) {
return 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(event_based_actor* self) {
// create two matrices with ascending values
matrix_type m1;
m1.iota_fill();
auto m2 = m1;
// print "source" matrix
cout << "calculating square of matrix:" << endl
<< to_string(m1) << 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 function operates on vectors,
// this function converts a tuple of two matrices
// to a tuple of vectors; note that this function returns
// an option (an empty results causes the actor to ignore
// the message)
[] (message msg) -> optional<cow_tuple<fvec, fvec>> {
auto opt = tuple_cast<matrix_type, matrix_type>(msg);
if (opt) {
return make_cow_tuple(move(get_ref<0>(*opt).data()),
move(get_ref<1>(*opt).data()));
}
return none;
},
// 4th arg: converts the ouptut vector back to a matrix that is then
// used as response message
[] (fvec& result) -> message {
return make_message(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
self->sync_send(worker, move(m1), move(m2)).then(
[](const matrix_type& result) {
cout << "result:" << endl << to_string(result);
}
);
}
int main() {
// matrix_type ist not a simple type,
// it must be annouced to libcaf
announce<matrix_type>();
spawn(multiplier);
await_all_actors_done();
shutdown();
return 0;
}
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2014 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and LICENCE_ALTERNATIVE. *
* *
* If you did not receive a copy of the license files, see *
* http://opensource.org/licenses/BSD-3-Clause and *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#include <vector>
#include <iomanip>
#include <numeric>
#include <iostream>
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
using namespace std;
using namespace caf;
namespace {
constexpr size_t matrix_size = 8;
constexpr const char* kernel_name = "matrix_mult";
// opencl kernel, multiplies matrix1 and matrix2
// last parameter is, by convention, the output parameter
constexpr const char* kernel_source = R"__(
__kernel void matrix_mult(__global float* matrix1,
__global float* matrix2,
__global float* output) {
// we only use square matrices, hence: width == height
size_t size = get_global_size(0); // == get_global_size_(1);
size_t x = get_global_id(0);
size_t y = get_global_id(1);
float result = 0;
for (size_t idx = 0; idx < size; ++idx) {
result += matrix1[idx + y * size] * matrix2[x + idx * size];
}
output[x+y*size] = result;
}
)__";
} // namespace <anonymous>
void print_as_matrix(const vector<float>& matrix) {
for (size_t column = 0; column < matrix_size; ++column) {
for (size_t row = 0; row < matrix_size; ++row) {
cout << fixed << setprecision(2) << setw(9)
<< matrix[row + column * matrix_size];
}
cout << endl;
}
}
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);
vector<float> m2(matrix_size * matrix_size);
// fill each with ascending values
iota(m1.begin(), m1.end(), 0);
iota(m2.begin(), m2.end(), 0);
// print "source" matrix
cout << "calculating square of matrix:" << endl;
print_as_matrix(m1);
cout << endl;
// spawn an opencl actor
// template parameter: signature of opencl kernel using proper return type
// instead of output parameter (implicitly
// mapped to the last kernel argument)
// 1st arg: source 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
// 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);
// send both matrices to the actor and wait for a result
self->sync_send(worker, move(m1), move(m2)).then(
[](const vector<float>& result) {
cout << "result: " << endl;
print_as_matrix(result);
}
);
}
int main() {
announce<vector<float>>();
spawn(multiplier);
await_all_actors_done();
shutdown();
return 0;
}
...@@ -37,7 +37,6 @@ const std::vector<device_info> opencl_metainfo::get_devices() const { ...@@ -37,7 +37,6 @@ const std::vector<device_info> opencl_metainfo::get_devices() const {
void opencl_metainfo::initialize() { void opencl_metainfo::initialize() {
cl_int err{0}; cl_int err{0};
// get number of available platforms // get number of available platforms
cl_uint number_of_platforms; cl_uint number_of_platforms;
err = clGetPlatformIDs(0, nullptr, &number_of_platforms); err = clGetPlatformIDs(0, nullptr, &number_of_platforms);
...@@ -48,7 +47,6 @@ void opencl_metainfo::initialize() { ...@@ -48,7 +47,6 @@ void opencl_metainfo::initialize() {
CAF_LOGMF(CAF_ERROR, oss.str()); CAF_LOGMF(CAF_ERROR, oss.str());
throw logic_error(oss.str()); throw logic_error(oss.str());
} }
// get platform ids // get platform ids
vector<cl_platform_id> ids(number_of_platforms); vector<cl_platform_id> ids(number_of_platforms);
err = clGetPlatformIDs(ids.size(), ids.data(), nullptr); err = clGetPlatformIDs(ids.size(), ids.data(), nullptr);
...@@ -58,7 +56,6 @@ void opencl_metainfo::initialize() { ...@@ -58,7 +56,6 @@ void opencl_metainfo::initialize() {
CAF_LOGMF(CAF_ERROR, oss.str()); CAF_LOGMF(CAF_ERROR, oss.str());
throw logic_error(oss.str()); throw logic_error(oss.str());
} }
// find gpu devices on our platform // find gpu devices on our platform
int pid{0}; int pid{0};
cl_uint num_devices{0}; cl_uint num_devices{0};
...@@ -85,14 +82,12 @@ void opencl_metainfo::initialize() { ...@@ -85,14 +82,12 @@ void opencl_metainfo::initialize() {
CAF_LOGMF(CAF_ERROR, oss.str()); CAF_LOGMF(CAF_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
auto pfn_notify = [](const char* errinfo, const void*, size_t, void*) { auto pfn_notify = [](const char* errinfo, const void*, size_t, void*) {
CAF_LOGC_ERROR("caf::opencl::opencl_metainfo", "initialize", CAF_LOGC_ERROR("caf::opencl::opencl_metainfo", "initialize",
"\n##### Error message via pfn_notify #####\n" + "\n##### Error message via pfn_notify #####\n" +
string(errinfo) + string(errinfo) +
"\n########################################"); "\n########################################");
}; };
// create a context // create a context
m_context.adopt(clCreateContext(0, devices.size(), devices.data(), pfn_notify, m_context.adopt(clCreateContext(0, devices.size(), devices.data(), pfn_notify,
nullptr, &err)); nullptr, &err));
...@@ -102,7 +97,6 @@ void opencl_metainfo::initialize() { ...@@ -102,7 +97,6 @@ void opencl_metainfo::initialize() {
CAF_LOGMF(CAF_ERROR, oss.str()); CAF_LOGMF(CAF_ERROR, oss.str());
throw runtime_error(oss.str()); throw runtime_error(oss.str());
} }
for (auto& d : devices) { for (auto& d : devices) {
CAF_LOG_TRACE("Creating command queue for device(s)."); CAF_LOG_TRACE("Creating command queue for device(s).");
device_ptr device; device_ptr device;
...@@ -161,7 +155,6 @@ void opencl_metainfo::initialize() { ...@@ -161,7 +155,6 @@ void opencl_metainfo::initialize() {
m_devices.push_back(move(dev_info)); m_devices.push_back(move(dev_info));
} }
} }
if (m_devices.empty()) { if (m_devices.empty()) {
ostringstream oss; ostringstream oss;
oss << "Could not create a command queue for " oss << "Could not create a command queue for "
...@@ -171,8 +164,13 @@ void opencl_metainfo::initialize() { ...@@ -171,8 +164,13 @@ void opencl_metainfo::initialize() {
} }
} }
void opencl_metainfo::dispose() { delete this; } void opencl_metainfo::dispose() {
delete this;
}
void opencl_metainfo::stop() {
// noop
}
} // namespace opencl } // namespace opencl
} // namespace caf } // namespace caf
cmake_minimum_required(VERSION 2.8) cmake_minimum_required(VERSION 2.8)
project(caf_unit_tests_opencl CXX) project(caf_unit_tests_opencl CXX)
add_custom_target(opencl_unit_tests) if (OPENCL_LIBRARIES)
add_custom_target(opencl_unit_tests)
include_directories(${LIBCAF_INCLUDE_DIRS}) include_directories(${LIBCAF_INCLUDE_DIRS})
macro(add_unit_test name)
macro(add_unit_test name) add_executable(test_${name} test_${name}.cpp ../../unit_testing/test.cpp ${ARGN})
add_executable(test_${name} test_${name}.cpp ../../unit_testing/test.cpp ${ARGN}) target_link_libraries(test_${name}
target_link_libraries(test_${name} ${LD_FLAGS}
${LD_FLAGS} ${LIBCAF_LIBRARIES}
${LIBCAF_LIBRARIES} ${PTHREAD_LIBRARIES}
${PTHREAD_LIBRARIES}) ${OPENCL_LIBRARIES})
add_test(${name} ${EXECUTABLE_OUTPUT_PATH}/test_${name}) add_test(${name} ${EXECUTABLE_OUTPUT_PATH}/test_${name})
add_dependencies(test_${name} all_unit_tests) add_dependencies(test_${name} all_unit_tests)
endmacro() endmacro()
add_unit_test(opencl)
add_unit_test(opencl) endif()
...@@ -4,7 +4,7 @@ ...@@ -4,7 +4,7 @@
#include <iostream> #include <iostream>
#include <algorithm> #include <algorithm>
#include "../../test.hpp" #include "../../unit_testing/test.hpp"
#include "caf/all.hpp" #include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp" #include "caf/opencl/spawn_cl.hpp"
...@@ -164,7 +164,7 @@ using matrix_type = square_matrix<matrix_size>; ...@@ -164,7 +164,7 @@ using matrix_type = square_matrix<matrix_size>;
size_t get_max_workgroup_size(size_t device_id, size_t dimension) { size_t get_max_workgroup_size(size_t device_id, size_t dimension) {
size_t max_size = 512; size_t max_size = 512;
auto devices = get_opencl_metainfo()->get_devices()[device_id]; auto devices = opencl_metainfo::instance()->get_devices()[device_id];
size_t dimsize = devices.get_max_work_items_per_dim()[dimension]; size_t dimsize = devices.get_max_work_items_per_dim()[dimension];
return max_size < dimsize ? max_size : dimsize; return max_size < dimsize ? max_size : dimsize;
} }
......
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