Skip to content
Projects
Groups
Snippets
Help
Loading...
Help
Support
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
A
Actor Framework
Project overview
Project overview
Details
Activity
Releases
Repository
Repository
Files
Commits
Branches
Tags
Contributors
Graph
Compare
Issues
0
Issues
0
List
Boards
Labels
Milestones
Merge Requests
0
Merge Requests
0
CI / CD
CI / CD
Pipelines
Jobs
Schedules
Operations
Operations
Metrics
Environments
Analytics
Analytics
CI / CD
Repository
Value Stream
Wiki
Wiki
Snippets
Snippets
Members
Members
Collapse sidebar
Close sidebar
Activity
Graph
Create a new issue
Jobs
Commits
Issue Boards
Open sidebar
cpp-libs
Actor Framework
Commits
f550bfe8
Commit
f550bfe8
authored
Nov 11, 2019
by
Dominik Charousset
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
Remove OpenCL module
parent
b05a4b7b
Changes
32
Hide whitespace changes
Inline
Side-by-side
Showing
32 changed files
with
2 additions
and
5272 deletions
+2
-5272
CMakeLists.txt
CMakeLists.txt
+0
-25
configure
configure
+2
-6
doc/CMakeLists.txt
doc/CMakeLists.txt
+0
-1
doc/tex/OpenCL.tex
doc/tex/OpenCL.tex
+0
-25
libcaf_opencl/CMakeLists.txt
libcaf_opencl/CMakeLists.txt
+0
-55
libcaf_opencl/README.md
libcaf_opencl/README.md
+0
-4
libcaf_opencl/caf/detail/command_helper.hpp
libcaf_opencl/caf/detail/command_helper.hpp
+0
-57
libcaf_opencl/caf/detail/raw_ptr.hpp
libcaf_opencl/caf/detail/raw_ptr.hpp
+0
-56
libcaf_opencl/caf/detail/spawn_helper.hpp
libcaf_opencl/caf/detail/spawn_helper.hpp
+0
-65
libcaf_opencl/caf/opencl/actor_facade.hpp
libcaf_opencl/caf/opencl/actor_facade.hpp
+0
-468
libcaf_opencl/caf/opencl/all.hpp
libcaf_opencl/caf/opencl/all.hpp
+0
-23
libcaf_opencl/caf/opencl/arguments.hpp
libcaf_opencl/caf/opencl/arguments.hpp
+0
-492
libcaf_opencl/caf/opencl/command.hpp
libcaf_opencl/caf/opencl/command.hpp
+0
-257
libcaf_opencl/caf/opencl/device.hpp
libcaf_opencl/caf/opencl/device.hpp
+0
-331
libcaf_opencl/caf/opencl/global.hpp
libcaf_opencl/caf/opencl/global.hpp
+0
-72
libcaf_opencl/caf/opencl/manager.hpp
libcaf_opencl/caf/opencl/manager.hpp
+0
-278
libcaf_opencl/caf/opencl/mem_ref.hpp
libcaf_opencl/caf/opencl/mem_ref.hpp
+0
-184
libcaf_opencl/caf/opencl/nd_range.hpp
libcaf_opencl/caf/opencl/nd_range.hpp
+0
-72
libcaf_opencl/caf/opencl/opencl_err.hpp
libcaf_opencl/caf/opencl/opencl_err.hpp
+0
-85
libcaf_opencl/caf/opencl/platform.hpp
libcaf_opencl/caf/opencl/platform.hpp
+0
-83
libcaf_opencl/caf/opencl/program.hpp
libcaf_opencl/caf/opencl/program.hpp
+0
-61
libcaf_opencl/examples/CMakeLists.txt
libcaf_opencl/examples/CMakeLists.txt
+0
-25
libcaf_opencl/examples/proper_matrix.cpp
libcaf_opencl/examples/proper_matrix.cpp
+0
-202
libcaf_opencl/examples/scan.cpp
libcaf_opencl/examples/scan.cpp
+0
-283
libcaf_opencl/examples/simple_matrix.cpp
libcaf_opencl/examples/simple_matrix.cpp
+0
-118
libcaf_opencl/src/device.cpp
libcaf_opencl/src/device.cpp
+0
-135
libcaf_opencl/src/global.cpp
libcaf_opencl/src/global.cpp
+0
-292
libcaf_opencl/src/manager.cpp
libcaf_opencl/src/manager.cpp
+0
-207
libcaf_opencl/src/opencl_err.cpp
libcaf_opencl/src/opencl_err.cpp
+0
-41
libcaf_opencl/src/platform.cpp
libcaf_opencl/src/platform.cpp
+0
-105
libcaf_opencl/src/program.cpp
libcaf_opencl/src/program.cpp
+0
-50
libcaf_opencl/test/opencl.cpp
libcaf_opencl/test/opencl.cpp
+0
-1114
No files found.
CMakeLists.txt
View file @
f550bfe8
...
...
@@ -138,7 +138,6 @@ pretty_no("CAF_NO_MEM_MANAGEMENT")
pretty_no
(
"CAF_NO_EXCEPTIONS"
)
pretty_no
(
"CAF_BUILD_STATIC_ONLY"
)
pretty_no
(
"CAF_BUILD_STATIC"
)
pretty_no
(
"CAF_NO_OPENCL"
)
pretty_no
(
"CAF_NO_OPENSSL"
)
pretty_no
(
"CAF_NO_PYTHON"
)
pretty_no
(
"CAF_NO_TOOLS"
)
...
...
@@ -451,11 +450,6 @@ set(LIBCAF_INCLUDE_DIRS
"
${
CMAKE_CURRENT_SOURCE_DIR
}
/libcaf_core"
"
${
CMAKE_CURRENT_SOURCE_DIR
}
/libcaf_io"
"
${
CMAKE_CURRENT_SOURCE_DIR
}
/libcaf_test"
)
# path to caf opencl headers
if
(
EXISTS
"
${
CMAKE_CURRENT_SOURCE_DIR
}
/libcaf_opencl/CMakeLists.txt"
)
set
(
LIBCAF_INCLUDE_DIRS
"
${
CMAKE_CURRENT_SOURCE_DIR
}
/libcaf_opencl/"
"
${
LIBCAF_INCLUDE_DIRS
}
"
)
endif
()
# enable tests if not disabled
if
(
NOT CAF_NO_UNIT_TESTS
)
enable_testing
()
...
...
@@ -576,23 +570,6 @@ if(NOT CAF_NO_OPENSSL)
endif
(
OPENSSL_FOUND
)
endif
()
# build opencl library if not told otherwise and OpenCL package was found
if
(
NOT CAF_NO_OPENCL
)
if
(
${
CMAKE_MAJOR_VERSION
}
.
${
CMAKE_MINOR_VERSION
}
GREATER 3.0
)
find_package
(
OpenCL
)
if
(
OpenCL_FOUND
)
include_directories
(
BEFORE
${
OpenCL_INCLUDE_DIRS
}
)
add_optional_caf_lib
(
opencl
)
add_optional_caf_binaries
(
libcaf_opencl/examples
)
else
()
set
(
CAF_NO_OPENCL yes
)
endif
()
else
()
set
(
CAF_NO_OPENCL yes
)
message
(
STATUS
"Could NOT find OpenCL, requires Cmake >= 3.1."
)
endif
()
endif
()
# build Python binding if not being told otherwise
if
(
NOT CAF_NO_PYTHON AND EXISTS
"
${
CMAKE_CURRENT_SOURCE_DIR
}
/libcaf_python/CMakeLists.txt"
)
add_subdirectory
(
libcaf_python
)
...
...
@@ -687,7 +664,6 @@ invertYesNo(CAF_NO_TOOLS CAF_BUILD_TOOLS)
invertYesNo
(
CAF_NO_UNIT_TESTS CAF_BUILD_UNIT_TESTS
)
invertYesNo
(
CAF_NO_EXCEPTIONS CAF_BUILD_WITH_EXCEPTIONS
)
invertYesNo
(
CAF_NO_MEM_MANAGEMENT CAF_BUILD_MEM_MANAGEMENT
)
invertYesNo
(
CAF_NO_OPENCL CAF_BUILD_OPENCL
)
invertYesNo
(
CAF_NO_OPENSSL CAF_BUILD_OPENSSL
)
invertYesNo
(
CAF_NO_PYTHON CAF_BUILD_PYTHON
)
# collect all compiler flags
...
...
@@ -715,7 +691,6 @@ if(NOT CAF_NO_SUMMARY)
"
\n
Build tools:
${
CAF_BUILD_TOOLS
}
"
"
\n
Build examples:
${
CAF_BUILD_EXAMPLES
}
"
"
\n
Build unit tests:
${
CAF_BUILD_UNIT_TESTS
}
"
"
\n
Build OpenCL:
${
CAF_BUILD_OPENCL
}
"
"
\n
Build OpenSSL:
${
CAF_BUILD_OPENSSL
}
"
"
\n
Build Python:
${
CAF_BUILD_PYTHON
}
"
"
\n
"
...
...
configure
View file @
f550bfe8
...
...
@@ -58,7 +58,6 @@ Usage: $0 [OPTION]... [VAR=VALUE]...
--no-examples build without examples
--no-curl-examples build without libcurl examples
--no-unit-tests build without unit tests
--no-opencl build without OpenCL module
--no-openssl build without OpenSSL module
--no-tools build without CAF tools such as caf-run
--no-io build without I/O module
...
...
@@ -66,7 +65,7 @@ Usage: $0 [OPTION]... [VAR=VALUE]...
--no-summary do not print configuration before building
--libs-only sets no-examples, no-tools, no-python,
and no-unit-tests
--core-only same as libs-only but also adds sets
no-opencl,
--core-only same as libs-only but also adds sets
no-openssl, and no-io
Debugging:
...
...
@@ -337,9 +336,6 @@ while [ $# -ne 0 ]; do
--no-unit-tests
)
append_cache_entry CAF_NO_UNIT_TESTS BOOL
yes
;;
--no-opencl
)
append_cache_entry CAF_NO_OPENCL BOOL
yes
;;
--no-openssl
)
append_cache_entry CAF_NO_OPENSSL BOOL
yes
;;
...
...
@@ -373,7 +369,7 @@ while [ $# -ne 0 ]; do
done
;;
--core-only
)
for
var
in
CAF_NO_TOOLS CAF_NO_PYTHON CAF_NO_EXAMPLES CAF_NO_UNIT_TESTS CAF_NO_IO CAF_NO_OPEN
CL CAF_NO_OPEN
SSL
;
do
for
var
in
CAF_NO_TOOLS CAF_NO_PYTHON CAF_NO_EXAMPLES CAF_NO_UNIT_TESTS CAF_NO_IO CAF_NO_OPENSSL
;
do
append_cache_entry
$var
BOOL
yes
done
;;
...
...
doc/CMakeLists.txt
View file @
f550bfe8
...
...
@@ -18,7 +18,6 @@ set(sources
tex/Messages.tex
tex/MigrationGuides.tex
tex/NetworkTransparency.tex
tex/OpenCL.tex
tex/ReferenceCounting.tex
tex/Registry.tex
tex/RemoteSpawn.tex
...
...
doc/tex/OpenCL.tex
deleted
100644 → 0
View file @
b05a4b7b
\section
{
OpenCL-based Actors
}
CAF supports transparent integration of OpenCL functions.
\begin{lstlisting}
// opencl kernel for matrix multiplication;
// 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;
}
)
__
";
\end{lstlisting}
libcaf_opencl/CMakeLists.txt
deleted
100644 → 0
View file @
b05a4b7b
project
(
caf_opencl C CXX
)
# get header files; only needed by CMake generators,
# e.g., for creating proper Xcode projects
file
(
GLOB_RECURSE LIBCAF_OPENCL_HDRS
"caf/*.hpp"
)
# list cpp files excluding platform-dependent files
set
(
LIBCAF_OPENCL_SRCS
src/device.cpp
src/global.cpp
src/manager.cpp
src/opencl_err.cpp
src/platform.cpp
src/program.cpp
)
# build shared library if not compiling static only
if
(
NOT CAF_BUILD_STATIC_ONLY
)
add_library
(
libcaf_opencl_shared SHARED
${
LIBCAF_OPENCL_SRCS
}
${
LIBCAF_OPENCL_HDRS
}
${
OpenCL_INCLUDE_DIRS
}
)
target_link_libraries
(
libcaf_opencl_shared
${
CAF_EXTRA_LDFLAGS
}
${
CAF_LIBRARY_CORE
}
${
OpenCL_LIBRARIES
}
)
target_link_libraries
(
libcaf_opencl_shared
$<BUILD_INTERFACE:
${
CMAKE_CURRENT_SOURCE_DIR
}
>
$<INSTALL_INTERFACE:include>
)
set_target_properties
(
libcaf_opencl_shared
PROPERTIES
SOVERSION
"
${
CAF_VERSION
}
"
VERSION
"
${
CAF_LIB_VERSION
}
"
OUTPUT_NAME caf_opencl
)
install
(
TARGETS libcaf_opencl_shared LIBRARY DESTINATION
${
CMAKE_INSTALL_LIBDIR
}
)
endif
()
# build static library only if --build-static or --build-static-only was set
if
(
CAF_BUILD_STATIC_ONLY OR CAF_BUILD_STATIC
)
add_library
(
libcaf_opencl_static STATIC
${
LIBCAF_OPENCL_SRCS
}
${
LIBCAF_OPENCL_HDRS
}
${
OpenCL_INCLUDE_DIRS
}
)
target_link_libraries
(
libcaf_opencl_static
${
CAF_EXTRA_LDFLAGS
}
${
CAF_LIBRARY_CORE_STATIC
}
${
OpenCL_LIBRARIES
}
)
target_link_libraries
(
libcaf_opencl_static
$<BUILD_INTERFACE:
${
CMAKE_CURRENT_SOURCE_DIR
}
>
$<INSTALL_INTERFACE:include>
)
set_target_properties
(
libcaf_opencl_static PROPERTIES OUTPUT_NAME caf_opencl_static
)
install
(
TARGETS libcaf_opencl_static ARCHIVE DESTINATION
${
CMAKE_INSTALL_LIBDIR
}
)
endif
()
# install includes
install
(
DIRECTORY
${
CMAKE_CURRENT_SOURCE_DIR
}
/caf
DESTINATION
${
CMAKE_INSTALL_INCLUDEDIR
}
FILES_MATCHING PATTERN
"*.hpp"
)
libcaf_opencl/README.md
deleted
100644 → 0
View file @
b05a4b7b
OpenCL Actors
=============
This module eases the use of OpenCL with CAF. See our
[
Wiki page
](
https://github.com/actor-framework/actor-framework/wiki/OpenCL-Actors
)
for details.
libcaf_opencl/caf/detail/command_helper.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include "caf/detail/type_list.hpp"
#include "caf/opencl/command.hpp"
namespace
caf
{
namespace
detail
{
// signature for the function that is applied to output arguments
template
<
class
List
>
struct
output_function_sig
;
template
<
class
...
Ts
>
struct
output_function_sig
<
detail
::
type_list
<
Ts
...
>>
{
using
type
=
std
::
function
<
message
(
Ts
&
...)
>
;
};
// derive signature of the command that handles the kernel execution
template
<
class
T
,
class
List
>
struct
command_sig
;
template
<
class
T
,
class
...
Ts
>
struct
command_sig
<
T
,
detail
::
type_list
<
Ts
...
>>
{
using
type
=
opencl
::
command
<
T
,
Ts
...
>
;
};
// derive type for a tuple matching the arguments as mem_refs
template
<
class
List
>
struct
tuple_type_of
;
template
<
class
...
Ts
>
struct
tuple_type_of
<
detail
::
type_list
<
Ts
...
>>
{
using
type
=
std
::
tuple
<
Ts
...
>
;
};
}
// namespace detail
}
// namespace caf
libcaf_opencl/caf/detail/raw_ptr.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <memory>
#include <algorithm>
#include <type_traits>
#include "caf/intrusive_ptr.hpp"
#include "caf/opencl/global.hpp"
#define CAF_OPENCL_PTR_ALIAS(aliasname, cltype, claddref, clrelease) \
inline void intrusive_ptr_add_ref(cltype ptr) { claddref(ptr); } \
inline void intrusive_ptr_release(cltype ptr) { clrelease(ptr); } \
namespace caf { \
namespace detail { \
using aliasname = intrusive_ptr<std::remove_pointer<cltype>::type>; \
}
/* namespace detail */
\
} // namespace caf
CAF_OPENCL_PTR_ALIAS
(
raw_mem_ptr
,
cl_mem
,
clRetainMemObject
,
clReleaseMemObject
)
CAF_OPENCL_PTR_ALIAS
(
raw_event_ptr
,
cl_event
,
clRetainEvent
,
clReleaseEvent
)
CAF_OPENCL_PTR_ALIAS
(
raw_kernel_ptr
,
cl_kernel
,
clRetainKernel
,
clReleaseKernel
)
CAF_OPENCL_PTR_ALIAS
(
raw_context_ptr
,
cl_context
,
clRetainContext
,
clReleaseContext
)
CAF_OPENCL_PTR_ALIAS
(
raw_program_ptr
,
cl_program
,
clRetainProgram
,
clReleaseProgram
)
CAF_OPENCL_PTR_ALIAS
(
raw_device_ptr
,
cl_device_id
,
clRetainDeviceDummy
,
clReleaseDeviceDummy
)
CAF_OPENCL_PTR_ALIAS
(
raw_command_queue_ptr
,
cl_command_queue
,
clRetainCommandQueue
,
clReleaseCommandQueue
)
libcaf_opencl/caf/detail/spawn_helper.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include "caf/opencl/actor_facade.hpp"
namespace
caf
{
namespace
detail
{
struct
tuple_construct
{
};
template
<
bool
PassConfig
,
class
...
Ts
>
struct
cl_spawn_helper
{
using
impl
=
opencl
::
actor_facade
<
PassConfig
,
Ts
...
>
;
using
map_in_fun
=
typename
impl
::
input_mapping
;
using
map_out_fun
=
typename
impl
::
output_mapping
;
actor
operator
()(
actor_config
actor_cfg
,
const
opencl
::
program_ptr
p
,
const
char
*
fn
,
const
opencl
::
nd_range
&
range
,
Ts
&&
...
xs
)
const
{
return
actor_cast
<
actor
>
(
impl
::
create
(
std
::
move
(
actor_cfg
),
p
,
fn
,
range
,
map_in_fun
{},
map_out_fun
{},
std
::
forward
<
Ts
>
(
xs
)...));
}
actor
operator
()(
actor_config
actor_cfg
,
const
opencl
::
program_ptr
p
,
const
char
*
fn
,
const
opencl
::
nd_range
&
range
,
map_in_fun
map_input
,
Ts
&&
...
xs
)
const
{
return
actor_cast
<
actor
>
(
impl
::
create
(
std
::
move
(
actor_cfg
),
p
,
fn
,
range
,
std
::
move
(
map_input
),
map_out_fun
{},
std
::
forward
<
Ts
>
(
xs
)...));
}
actor
operator
()(
actor_config
actor_cfg
,
const
opencl
::
program_ptr
p
,
const
char
*
fn
,
const
opencl
::
nd_range
&
range
,
map_in_fun
map_input
,
map_out_fun
map_output
,
Ts
&&
...
xs
)
const
{
return
actor_cast
<
actor
>
(
impl
::
create
(
std
::
move
(
actor_cfg
),
p
,
fn
,
range
,
std
::
move
(
map_input
),
std
::
move
(
map_output
),
std
::
forward
<
Ts
>
(
xs
)...));
}
};
}
// namespace detail
}
// namespace caf
libcaf_opencl/caf/opencl/actor_facade.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <ostream>
#include <iostream>
#include <algorithm>
#include <stdexcept>
#include "caf/all.hpp"
#include "caf/intrusive_ptr.hpp"
#include "caf/raise_error.hpp"
#include "caf/detail/raw_ptr.hpp"
#include "caf/detail/command_helper.hpp"
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/command.hpp"
#include "caf/opencl/mem_ref.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/nd_range.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/opencl_err.hpp"
namespace
caf
{
namespace
opencl
{
class
manager
;
template
<
bool
PassConfig
,
class
...
Ts
>
class
actor_facade
:
public
local_actor
{
public:
using
arg_types
=
detail
::
type_list
<
Ts
...
>
;
using
unpacked_types
=
typename
detail
::
tl_map
<
arg_types
,
extract_type
>::
type
;
using
input_wrapped_types
=
typename
detail
::
tl_filter
<
arg_types
,
is_input_arg
>::
type
;
using
input_types
=
typename
detail
::
tl_map
<
input_wrapped_types
,
extract_input_type
>::
type
;
using
input_mapping
=
typename
std
::
conditional
<
PassConfig
,
std
::
function
<
optional
<
message
>
(
nd_range
&
,
message
&
)
>
,
std
::
function
<
optional
<
message
>
(
message
&
)
>
>::
type
;
using
output_wrapped_types
=
typename
detail
::
tl_filter
<
arg_types
,
is_output_arg
>::
type
;
using
output_types
=
typename
detail
::
tl_map
<
output_wrapped_types
,
extract_output_type
>::
type
;
using
output_mapping
=
typename
detail
::
output_function_sig
<
output_types
>::
type
;
using
processing_list
=
typename
cl_arg_info_list
<
arg_types
>::
type
;
using
command_type
=
typename
detail
::
command_sig
<
actor_facade
,
output_types
>::
type
;
typename
detail
::
il_indices
<
arg_types
>::
type
indices
;
using
evnt_vec
=
std
::
vector
<
cl_event
>
;
using
mem_vec
=
std
::
vector
<
detail
::
raw_mem_ptr
>
;
using
len_vec
=
std
::
vector
<
size_t
>
;
using
out_tup
=
typename
detail
::
tuple_type_of
<
output_types
>::
type
;
const
char
*
name
()
const
override
{
return
"OpenCL actor"
;
}
static
actor
create
(
actor_config
actor_conf
,
const
program_ptr
prog
,
const
char
*
kernel_name
,
const
nd_range
&
range
,
input_mapping
map_args
,
output_mapping
map_result
,
Ts
&&
...
xs
)
{
if
(
range
.
dimensions
().
empty
())
CAF_RAISE_ERROR
(
"OpenCL kernel needs at least 1 global dimension"
);
auto
check_vec
=
[
&
](
const
dim_vec
&
vec
)
{
if
(
!
vec
.
empty
()
&&
vec
.
size
()
!=
range
.
dimensions
().
size
())
CAF_RAISE_ERROR
(
"illegal vector size"
);
};
check_vec
(
range
.
offsets
());
check_vec
(
range
.
local_dimensions
());
auto
&
sys
=
actor_conf
.
host
->
system
();
auto
itr
=
prog
->
available_kernels_
.
find
(
kernel_name
);
if
(
itr
==
prog
->
available_kernels_
.
end
())
{
detail
::
raw_kernel_ptr
kernel
;
kernel
.
reset
(
v2get
(
CAF_CLF
(
clCreateKernel
),
prog
->
program_
.
get
(),
kernel_name
),
false
);
return
make_actor
<
actor_facade
,
actor
>
(
sys
.
next_actor_id
(),
sys
.
node
(),
&
sys
,
std
::
move
(
actor_conf
),
prog
,
kernel
,
range
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward_as_tuple
(
xs
...));
}
return
make_actor
<
actor_facade
,
actor
>
(
sys
.
next_actor_id
(),
sys
.
node
(),
&
sys
,
std
::
move
(
actor_conf
),
prog
,
itr
->
second
,
range
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward_as_tuple
(
xs
...));
}
void
enqueue
(
strong_actor_ptr
sender
,
message_id
mid
,
message
content
,
response_promise
promise
)
{
CAF_PUSH_AID
(
id
());
CAF_LOG_TRACE
(
""
);
if
(
!
map_arguments
(
content
))
return
;
if
(
!
content
.
match_elements
(
input_types
{}))
{
CAF_LOG_ERROR
(
"Message types do not match the expected signature."
);
return
;
}
auto
hdl
=
std
::
make_tuple
(
sender
,
mid
.
response_id
());
evnt_vec
events
;
mem_vec
input_buffers
;
mem_vec
output_buffers
;
mem_vec
scratch_buffers
;
len_vec
result_lengths
;
out_tup
result
;
add_kernel_arguments
(
events
,
// accumulate events for execution
input_buffers
,
// opencl buffers included in in msg
output_buffers
,
// opencl buffers included in out msg
scratch_buffers
,
// opencl only used here
result
,
// tuple to save the output values
result_lengths
,
// size of buffers to read back
content
,
// message content
indices
);
// enable extraction of types from msg
auto
cmd
=
make_counted
<
command_type
>
(
std
::
move
(
promise
),
actor_cast
<
strong_actor_ptr
>
(
this
),
std
::
move
(
events
),
std
::
move
(
input_buffers
),
std
::
move
(
output_buffers
),
std
::
move
(
scratch_buffers
),
std
::
move
(
result_lengths
),
std
::
move
(
content
),
std
::
move
(
result
),
range_
);
cmd
->
enqueue
();
}
void
enqueue
(
mailbox_element_ptr
ptr
,
execution_unit
*
)
override
{
CAF_ASSERT
(
ptr
!=
nullptr
);
CAF_LOG_TRACE
(
CAF_ARG
(
*
ptr
));
response_promise
promise
{
ctrl
(),
*
ptr
};
enqueue
(
ptr
->
sender
,
ptr
->
mid
,
ptr
->
move_content_to_message
(),
std
::
move
(
promise
));
}
void
enqueue
(
strong_actor_ptr
sender
,
message_id
mid
,
message
content
,
execution_unit
*
host
)
override
{
CAF_LOG_TRACE
(
""
);
enqueue
(
make_mailbox_element
(
std
::
move
(
sender
),
mid
,
{},
std
::
move
(
content
)),
host
);
}
actor_facade
(
actor_config
actor_conf
,
const
program_ptr
prog
,
detail
::
raw_kernel_ptr
kernel
,
nd_range
range
,
input_mapping
map_args
,
output_mapping
map_result
,
std
::
tuple
<
Ts
...
>
xs
)
:
local_actor
(
actor_conf
),
kernel_
(
std
::
move
(
kernel
)),
program_
(
prog
->
program_
),
context_
(
prog
->
context_
),
queue_
(
prog
->
queue_
),
range_
(
std
::
move
(
range
)),
map_args_
(
std
::
move
(
map_args
)),
map_results_
(
std
::
move
(
map_result
)),
kernel_signature_
(
std
::
move
(
xs
))
{
CAF_LOG_TRACE
(
CAF_ARG
(
this
->
id
()));
default_length_
=
std
::
accumulate
(
std
::
begin
(
range_
.
dimensions
()),
std
::
end
(
range_
.
dimensions
()),
size_t
{
1
},
std
::
multiplies
<
size_t
>
{});
}
void
add_kernel_arguments
(
evnt_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
,
len_vec
&
,
message
&
,
detail
::
int_list
<>
)
{
// nop
}
/// The separation into input, output and scratch buffers is required to
/// access the related memory handles later on. The scratch and input handles
/// are saved to prevent deletion before the kernel finished execution.
template
<
long
I
,
long
...
Is
>
void
add_kernel_arguments
(
evnt_vec
&
events
,
mem_vec
&
inputs
,
mem_vec
&
outputs
,
mem_vec
&
scratch
,
out_tup
&
result
,
len_vec
&
lengths
,
message
&
msg
,
detail
::
int_list
<
I
,
Is
...
>
)
{
using
arg_type
=
typename
detail
::
tl_at
<
processing_list
,
I
>::
type
;
create_buffer
<
I
,
arg_type
::
in_pos
,
arg_type
::
out_pos
>
(
std
::
get
<
I
>
(
kernel_signature_
),
events
,
lengths
,
inputs
,
outputs
,
scratch
,
result
,
msg
);
add_kernel_arguments
(
events
,
inputs
,
outputs
,
scratch
,
result
,
lengths
,
msg
,
detail
::
int_list
<
Is
...
>
{});
}
// Two functions to handle `in` arguments: val and mref
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
in
<
T
,
val
>&
,
evnt_vec
&
events
,
len_vec
&
,
mem_vec
&
inputs
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
using
container_type
=
std
::
vector
<
value_type
>
;
auto
&
container
=
msg
.
get_as
<
container_type
>
(
InPos
);
auto
len
=
container
.
size
();
size_t
num_bytes
=
sizeof
(
value_type
)
*
len
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
size_t
{
CL_MEM_READ_WRITE
},
num_bytes
,
nullptr
);
auto
event
=
v1get
<
cl_event
>
(
CAF_CLF
(
clEnqueueWriteBuffer
),
queue_
.
get
(),
buffer
,
0u
,
// --> CL_FALSE,
0u
,
num_bytes
,
container
.
data
());
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
buffer
));
events
.
push_back
(
event
);
inputs
.
emplace_back
(
buffer
,
false
);
}
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
in
<
T
,
mref
>&
,
evnt_vec
&
events
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
using
container_type
=
mem_ref
<
value_type
>
;
auto
container
=
msg
.
get_as
<
container_type
>
(
InPos
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
container
.
get
()));
auto
event
=
container
.
take_event
();
if
(
event
)
events
.
push_back
(
event
);
}
// Four functions to handle `in_out` arguments:
// val->val, val->mref, mref->val, mref->mref
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
in_out
<
T
,
val
,
val
>&
,
evnt_vec
&
events
,
len_vec
&
lengths
,
mem_vec
&
,
mem_vec
&
outputs
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
using
container_type
=
std
::
vector
<
value_type
>
;
auto
&
container
=
msg
.
get_as
<
container_type
>
(
InPos
);
auto
len
=
container
.
size
();
size_t
num_bytes
=
sizeof
(
value_type
)
*
len
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
size_t
{
CL_MEM_READ_WRITE
},
num_bytes
,
nullptr
);
auto
event
=
v1get
<
cl_event
>
(
CAF_CLF
(
clEnqueueWriteBuffer
),
queue_
.
get
(),
buffer
,
0u
,
// --> CL_FALSE,
0u
,
num_bytes
,
container
.
data
());
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
buffer
));
lengths
.
push_back
(
len
);
events
.
push_back
(
event
);
outputs
.
emplace_back
(
buffer
,
false
);
}
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
in_out
<
T
,
val
,
mref
>&
,
evnt_vec
&
events
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
result
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
using
container_type
=
std
::
vector
<
value_type
>
;
auto
&
container
=
msg
.
get_as
<
container_type
>
(
InPos
);
auto
len
=
container
.
size
();
size_t
num_bytes
=
sizeof
(
value_type
)
*
len
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
size_t
{
CL_MEM_READ_WRITE
},
num_bytes
,
nullptr
);
auto
event
=
v1get
<
cl_event
>
(
CAF_CLF
(
clEnqueueWriteBuffer
),
queue_
.
get
(),
buffer
,
0u
,
// --> CL_FALSE,
0u
,
num_bytes
,
container
.
data
());
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
buffer
));
events
.
push_back
(
event
);
std
::
get
<
OutPos
>
(
result
)
=
mem_ref
<
value_type
>
{
len
,
queue_
,
detail
::
raw_mem_ptr
{
buffer
,
false
},
size_t
{
CL_MEM_READ_WRITE
|
CL_MEM_HOST_READ_ONLY
},
nullptr
};
}
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
in_out
<
T
,
mref
,
val
>&
,
evnt_vec
&
events
,
len_vec
&
lengths
,
mem_vec
&
,
mem_vec
&
outputs
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
using
container_type
=
mem_ref
<
value_type
>
;
auto
container
=
msg
.
get_as
<
container_type
>
(
InPos
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
container
.
get
()));
auto
event
=
container
.
take_event
();
if
(
event
)
events
.
push_back
(
event
);
lengths
.
push_back
(
container
.
size
());
outputs
.
push_back
(
container
.
get
());
}
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
in_out
<
T
,
mref
,
mref
>&
,
evnt_vec
&
events
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
result
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
using
container_type
=
mem_ref
<
value_type
>
;
auto
container
=
msg
.
get_as
<
container_type
>
(
InPos
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
container
.
get
()));
auto
event
=
container
.
take_event
();
if
(
event
)
events
.
push_back
(
event
);
std
::
get
<
OutPos
>
(
result
)
=
container
;
}
// Two functions to handle `out` arguments: val and mref
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
out
<
T
,
val
>&
wrapper
,
evnt_vec
&
,
len_vec
&
lengths
,
mem_vec
&
,
mem_vec
&
outputs
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
auto
len
=
argument_length
(
wrapper
,
msg
,
default_length_
);
auto
num_bytes
=
sizeof
(
value_type
)
*
len
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
size_t
{
CL_MEM_READ_WRITE
|
CL_MEM_HOST_READ_ONLY
},
num_bytes
,
nullptr
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
buffer
));
outputs
.
emplace_back
(
buffer
,
false
);
lengths
.
push_back
(
len
);
}
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
out
<
T
,
mref
>&
wrapper
,
evnt_vec
&
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
result
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
auto
len
=
argument_length
(
wrapper
,
msg
,
default_length_
);
auto
num_bytes
=
sizeof
(
value_type
)
*
len
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
size_t
{
CL_MEM_READ_WRITE
|
CL_MEM_HOST_READ_ONLY
},
num_bytes
,
nullptr
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
buffer
));
std
::
get
<
OutPos
>
(
result
)
=
mem_ref
<
value_type
>
{
len
,
queue_
,
{
buffer
,
false
},
size_t
{
CL_MEM_READ_WRITE
|
CL_MEM_HOST_READ_ONLY
},
nullptr
};
}
// One function to handle `scratch` buffers
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
scratch
<
T
>&
wrapper
,
evnt_vec
&
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
scratch
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
auto
len
=
argument_length
(
wrapper
,
msg
,
default_length_
);
auto
num_bytes
=
sizeof
(
value_type
)
*
len
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
size_t
{
CL_MEM_READ_WRITE
|
CL_MEM_HOST_NO_ACCESS
},
num_bytes
,
nullptr
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
buffer
));
scratch
.
emplace_back
(
buffer
,
false
);
}
// One functions to handle `local` arguments
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
local
<
T
>&
wrapper
,
evnt_vec
&
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
auto
len
=
wrapper
(
msg
);
auto
num_bytes
=
sizeof
(
value_type
)
*
len
;
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
num_bytes
,
nullptr
);
}
// Two functions to handle `priv` arguments: val and hidden
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
priv
<
T
,
val
>&
,
evnt_vec
&
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
using
value_type
=
typename
detail
::
tl_at
<
unpacked_types
,
I
>::
type
;
auto
value_size
=
sizeof
(
value_type
);
auto
&
value
=
msg
.
get_as
<
value_type
>
(
InPos
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
value_size
,
static_cast
<
const
void
*>
(
&
value
));
}
template
<
long
I
,
int
InPos
,
int
OutPos
,
class
T
>
void
create_buffer
(
const
priv
<
T
,
hidden
>&
wrapper
,
evnt_vec
&
,
len_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
,
message
&
msg
)
{
auto
value_size
=
sizeof
(
T
);
auto
value
=
wrapper
(
msg
);
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
static_cast
<
unsigned
>
(
I
),
value_size
,
static_cast
<
const
void
*>
(
&
value
));
}
/// Helper function to calculate the elements in a buffer from in and out
/// argument wrappers.
template
<
class
Fun
>
size_t
argument_length
(
Fun
&
f
,
message
&
m
,
size_t
fallback
)
{
auto
length
=
f
(
m
);
return
length
&&
(
*
length
>
0
)
?
*
length
:
fallback
;
}
// Map function requires only the message as argument
template
<
bool
Q
=
PassConfig
>
detail
::
enable_if_t
<!
Q
,
bool
>
map_arguments
(
message
&
content
)
{
if
(
map_args_
)
{
auto
mapped
=
map_args_
(
content
);
if
(
!
mapped
)
{
CAF_LOG_ERROR
(
"Mapping argumentes failed."
);
return
false
;
}
content
=
std
::
move
(
*
mapped
);
}
return
true
;
}
// Map function requires reference to config as well as the message
template
<
bool
Q
=
PassConfig
>
detail
::
enable_if_t
<
Q
,
bool
>
map_arguments
(
message
&
content
)
{
if
(
map_args_
)
{
auto
mapped
=
map_args_
(
range_
,
content
);
if
(
!
mapped
)
{
CAF_LOG_ERROR
(
"Mapping argumentes failed."
);
return
false
;
}
content
=
std
::
move
(
*
mapped
);
}
return
true
;
}
void
launch
(
execution_unit
*
,
bool
,
bool
)
override
{
CAF_RAISE_ERROR
(
"launch of the actor facade should not be called"
);
}
detail
::
raw_kernel_ptr
kernel_
;
detail
::
raw_program_ptr
program_
;
detail
::
raw_context_ptr
context_
;
detail
::
raw_command_queue_ptr
queue_
;
nd_range
range_
;
input_mapping
map_args_
;
output_mapping
map_results_
;
std
::
tuple
<
Ts
...
>
kernel_signature_
;
size_t
default_length_
;
};
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/all.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Raphael Hiesgen <raphael.hiesgen (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 LICENSE_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. *
******************************************************************************/
#pragma once
#include "caf/opencl/manager.hpp"
libcaf_opencl/caf/opencl/arguments.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Raphael Hiesgen <raphael.hiesgen (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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <functional>
#include <type_traits>
#include "caf/message.hpp"
#include "caf/optional.hpp"
#include "caf/opencl/mem_ref.hpp"
namespace
caf
{
namespace
detail
{
template
<
class
T
,
class
F
>
std
::
function
<
optional
<
T
>
(
message
&
)
>
res_or_none
(
F
fun
)
{
return
[
fun
](
message
&
msg
)
->
optional
<
T
>
{
auto
res
=
msg
.
apply
(
fun
);
T
result
;
if
(
res
)
{
res
->
apply
([
&
](
size_t
x
)
{
result
=
x
;
});
return
result
;
}
return
none
;
};
}
template
<
class
F
,
class
T
>
T
try_apply_fun
(
F
&
fun
,
message
&
msg
,
const
T
&
fallback
)
{
if
(
fun
)
{
auto
res
=
fun
(
msg
);
if
(
res
)
return
*
res
;
}
return
fallback
;
}
}
// namespace detail
namespace
opencl
{
// Tag classes to mark arguments received in a messages as reference or value
/// Arguments tagged as `val` are expected as a vector (or value in case
/// of a private argument).
struct
val
{};
/// Arguments tagged as `mref` are expected as mem_ref, which is can be returned
/// by other opencl actors.
struct
mref
{};
/// Arguments tagged as `hidden` are created by the actor, using the config
/// passed in the argument wrapper. Only available for local and priv arguments.
struct
hidden
{};
/// Use as a default way to calculate output size. 0 will be set to the number
/// of work items at runtime.
struct
dummy_size_calculator
{
template
<
class
...
Ts
>
size_t
operator
()(
Ts
&&
...)
const
{
return
0
;
}
};
/// Common parent for opencl argument tags for the spawn function.
struct
arg_tag
{};
/// Empty tag as an alternative for conditional inheritance.
struct
empty_tag
{};
/// Tags the argument as input which requires initialization through a message.
struct
input_tag
{};
/// Tags the argument as output which includes its buffer in the result message.
struct
output_tag
{};
/// Tags the argument to require specification of the size of its buffer.
struct
requires_size_tag
{};
/// Tags the argument as a reference and not a value.
struct
is_ref_tag
;
/// Mark a spawn argument as input only
template
<
class
Arg
,
class
Tag
=
val
>
struct
in
:
arg_tag
,
input_tag
{
static_assert
(
std
::
is_same
<
Tag
,
val
>::
value
||
std
::
is_same
<
Tag
,
mref
>::
value
,
"Argument of type `in` must be passed as value or mem_ref."
);
using
tag_type
=
Tag
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
};
/// Mark a spawn argument as input and output
template
<
class
Arg
,
class
TagIn
=
val
,
class
TagOut
=
val
>
struct
in_out
:
arg_tag
,
input_tag
,
output_tag
{
static_assert
(
std
::
is_same
<
TagIn
,
val
>::
value
||
std
::
is_same
<
TagIn
,
mref
>::
value
,
"Argument of type `in_out` must be passed as value or mem_ref."
);
static_assert
(
std
::
is_same
<
TagOut
,
val
>::
value
||
std
::
is_same
<
TagOut
,
mref
>::
value
,
"Argument of type `in_out` must be returned as value or mem_ref."
);
using
tag_in_type
=
TagIn
;
using
tag_out_type
=
TagOut
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
};
/// Mark a spawn argument as output only
template
<
class
Arg
,
class
Tag
=
val
>
struct
out
:
arg_tag
,
output_tag
,
requires_size_tag
{
static_assert
(
std
::
is_same
<
Tag
,
val
>::
value
||
std
::
is_same
<
Tag
,
mref
>::
value
,
"Argument of type `out` must be returned as value or mem_ref."
);
using
tag_type
=
Tag
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
out
()
=
default
;
template
<
class
F
>
out
(
F
fun
)
:
fun_
{
detail
::
res_or_none
<
size_t
>
(
fun
)}
{
// nop
}
optional
<
size_t
>
operator
()(
message
&
msg
)
const
{
return
detail
::
try_apply_fun
(
fun_
,
msg
,
0UL
);
}
std
::
function
<
optional
<
size_t
>
(
message
&
)
>
fun_
;
};
/// Mark a spawn argument as on-device scratch space
template
<
class
Arg
>
struct
scratch
:
arg_tag
,
requires_size_tag
{
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
scratch
()
=
default
;
template
<
class
F
>
scratch
(
F
fun
)
:
fun_
{
detail
::
res_or_none
<
size_t
>
(
fun
)}
{
// nop
}
optional
<
size_t
>
operator
()(
message
&
msg
)
const
{
return
detail
::
try_apply_fun
(
fun_
,
msg
,
0UL
);
}
std
::
function
<
optional
<
size_t
>
(
message
&
)
>
fun_
;
};
/// Mark a spawn argument as a local memory argument. This argument cannot be
/// initalized from the CPU, but requires specification of its size. An
/// optional function allows calculation of the size depeding on the input
/// message.
template
<
class
Arg
>
struct
local
:
arg_tag
,
requires_size_tag
{
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
local
()
=
default
;
local
(
size_t
size
)
:
size_
(
size
)
{
}
template
<
class
F
>
local
(
size_t
size
,
F
fun
)
:
size_
(
size
),
fun_
{
detail
::
res_or_none
<
size_t
>
(
fun
)}
{
// nop
}
size_t
operator
()(
message
&
msg
)
const
{
return
detail
::
try_apply_fun
(
fun_
,
msg
,
size_
);
}
size_t
size_
;
std
::
function
<
optional
<
size_t
>
(
message
&
)
>
fun_
;
};
/// Mark a spawn argument as a private argument. Requires a default value but
/// can optionally be calculated depending on the input through a passed
/// function.
template
<
class
Arg
,
class
Tag
=
hidden
>
struct
priv
:
arg_tag
,
std
::
conditional
<
std
::
is_same
<
Tag
,
val
>::
value
,
input_tag
,
empty_tag
>::
type
{
static_assert
(
std
::
is_same
<
Tag
,
val
>::
value
||
std
::
is_same
<
Tag
,
hidden
>::
value
,
"Argument of type `priv` must be either a value or hidden."
);
using
tag_type
=
Tag
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
priv
()
=
default
;
priv
(
Arg
val
)
:
value_
(
val
)
{
static_assert
(
std
::
is_same
<
Tag
,
hidden
>::
value
,
"Argument of type `priv` can only be initialized with a value"
" if it is tagged as hidden."
);
}
template
<
class
F
>
priv
(
Arg
val
,
F
fun
)
:
value_
(
val
),
fun_
{
detail
::
res_or_none
<
Arg
>
(
fun
)}
{
static_assert
(
std
::
is_same
<
Tag
,
hidden
>::
value
,
"Argument of type `priv` can only be initialized with a value"
" if it is tagged as hidden."
);
}
Arg
operator
()(
message
&
msg
)
const
{
return
detail
::
try_apply_fun
(
fun_
,
msg
,
value_
);
}
Arg
value_
;
std
::
function
<
optional
<
Arg
>
(
message
&
)
>
fun_
;
};
///Cconverts C arrays, i.e., pointers, to vectors.
template
<
class
T
>
struct
carr_to_vec
{
using
type
=
T
;
};
template
<
class
T
>
struct
carr_to_vec
<
T
*>
{
using
type
=
std
::
vector
<
T
>
;
};
/// Filter types for any argument type.
template
<
class
T
>
struct
is_opencl_arg
:
std
::
is_base_of
<
arg_tag
,
T
>
{};
/// Filter type lists for input arguments
template
<
class
T
>
struct
is_input_arg
:
std
::
is_base_of
<
input_tag
,
T
>
{};
/// Filter type lists for output arguments
template
<
class
T
>
struct
is_output_arg
:
std
::
is_base_of
<
output_tag
,
T
>
{};
/// Filter for arguments that require size
template
<
class
T
>
struct
requires_size_arg
:
std
::
is_base_of
<
requires_size_tag
,
T
>
{};
/// Filter mem_refs
template
<
class
T
>
struct
is_ref_type
:
std
::
is_base_of
<
is_ref_tag
,
T
>
{};
template
<
class
T
>
struct
is_val_type
:
std
::
integral_constant
<
bool
,
!
std
::
is_base_of
<
is_ref_tag
,
T
>::
value
>
{};
/// extract types
template
<
class
T
>
struct
extract_type
{
};
template
<
class
T
,
class
Tag
>
struct
extract_type
<
in
<
T
,
Tag
>>
{
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
,
class
TagIn
,
class
TagOut
>
struct
extract_type
<
in_out
<
T
,
TagIn
,
TagOut
>>
{
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
,
class
Tag
>
struct
extract_type
<
out
<
T
,
Tag
>>
{
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
>
struct
extract_type
<
scratch
<
T
>>
{
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
>
struct
extract_type
<
local
<
T
>>
{
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
,
class
Tag
>
struct
extract_type
<
priv
<
T
,
Tag
>>
{
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
/// extract type expected in an incoming message
template
<
class
T
>
struct
extract_input_type
{
};
template
<
class
Arg
>
struct
extract_input_type
<
in
<
Arg
,
val
>>
{
using
type
=
std
::
vector
<
Arg
>
;
};
template
<
class
Arg
>
struct
extract_input_type
<
in
<
Arg
,
mref
>>
{
using
type
=
opencl
::
mem_ref
<
Arg
>
;
};
template
<
class
Arg
,
class
TagOut
>
struct
extract_input_type
<
in_out
<
Arg
,
val
,
TagOut
>>
{
using
type
=
std
::
vector
<
Arg
>
;
};
template
<
class
Arg
,
class
TagOut
>
struct
extract_input_type
<
in_out
<
Arg
,
mref
,
TagOut
>>
{
using
type
=
opencl
::
mem_ref
<
Arg
>
;
};
template
<
class
Arg
>
struct
extract_input_type
<
priv
<
Arg
,
val
>>
{
using
type
=
Arg
;
};
/// extract type sent in an outgoing message
template
<
class
T
>
struct
extract_output_type
{
};
template
<
class
Arg
>
struct
extract_output_type
<
out
<
Arg
,
val
>>
{
using
type
=
std
::
vector
<
Arg
>
;
};
template
<
class
Arg
>
struct
extract_output_type
<
out
<
Arg
,
mref
>>
{
using
type
=
opencl
::
mem_ref
<
Arg
>
;
};
template
<
class
Arg
,
class
TagIn
>
struct
extract_output_type
<
in_out
<
Arg
,
TagIn
,
val
>>
{
using
type
=
std
::
vector
<
Arg
>
;
};
template
<
class
Arg
,
class
TagIn
>
struct
extract_output_type
<
in_out
<
Arg
,
TagIn
,
mref
>>
{
using
type
=
opencl
::
mem_ref
<
Arg
>
;
};
/// extract input tag
template
<
class
T
>
struct
extract_input_tag
{
};
template
<
class
Arg
,
class
Tag
>
struct
extract_input_tag
<
in
<
Arg
,
Tag
>>
{
using
tag
=
Tag
;
};
template
<
class
Arg
,
class
TagIn
,
class
TagOut
>
struct
extract_input_tag
<
in_out
<
Arg
,
TagIn
,
TagOut
>>
{
using
tag
=
TagIn
;
};
template
<
class
Arg
>
struct
extract_input_tag
<
priv
<
Arg
,
val
>>
{
using
tag
=
val
;
};
/// extract output tag
template
<
class
T
>
struct
extract_output_tag
{
};
template
<
class
Arg
,
class
Tag
>
struct
extract_output_tag
<
out
<
Arg
,
Tag
>>
{
using
tag
=
Tag
;
};
template
<
class
Arg
,
class
TagIn
,
class
TagOut
>
struct
extract_output_tag
<
in_out
<
Arg
,
TagIn
,
TagOut
>>
{
using
tag
=
TagOut
;
};
/// Create the return message from tuple arumgent
struct
message_from_results
{
template
<
class
T
,
class
...
Ts
>
message
operator
()(
T
&
x
,
Ts
&
...
xs
)
{
return
make_message
(
std
::
move
(
x
),
std
::
move
(
xs
)...);
}
template
<
class
...
Ts
>
message
operator
()(
std
::
tuple
<
Ts
...
>&
values
)
{
return
apply_args
(
*
this
,
detail
::
get_indices
(
values
),
values
);
}
};
/// Calculate output indices from the kernel message
// index in output tuple
template
<
int
Counter
,
class
Arg
>
struct
out_index_of
{
static
constexpr
int
value
=
-
1
;
static
constexpr
int
next
=
Counter
;
};
template
<
int
Counter
,
class
Arg
,
class
TagIn
,
class
TagOut
>
struct
out_index_of
<
Counter
,
in_out
<
Arg
,
TagIn
,
TagOut
>>
{
static
constexpr
int
value
=
Counter
;
static
constexpr
int
next
=
Counter
+
1
;
};
template
<
int
Counter
,
class
Arg
,
class
Tag
>
struct
out_index_of
<
Counter
,
out
<
Arg
,
Tag
>>
{
static
constexpr
int
value
=
Counter
;
static
constexpr
int
next
=
Counter
+
1
;
};
// index in input message
template
<
int
Counter
,
class
Arg
>
struct
in_index_of
{
static
constexpr
int
value
=
-
1
;
static
constexpr
int
next
=
Counter
;
};
template
<
int
Counter
,
class
Arg
,
class
Tag
>
struct
in_index_of
<
Counter
,
in
<
Arg
,
Tag
>>
{
static
constexpr
int
value
=
Counter
;
static
constexpr
int
next
=
Counter
+
1
;
};
template
<
int
Counter
,
class
Arg
,
class
TagIn
,
class
TagOut
>
struct
in_index_of
<
Counter
,
in_out
<
Arg
,
TagIn
,
TagOut
>>
{
static
constexpr
int
value
=
Counter
;
static
constexpr
int
next
=
Counter
+
1
;
};
template
<
int
Counter
,
class
Arg
>
struct
in_index_of
<
Counter
,
priv
<
Arg
,
val
>>
{
static
constexpr
int
value
=
Counter
;
static
constexpr
int
next
=
Counter
+
1
;
};
template
<
int
In
,
int
Out
,
class
T
>
struct
cl_arg_info
{
static
constexpr
int
in_pos
=
In
;
static
constexpr
int
out_pos
=
Out
;
using
type
=
T
;
};
template
<
class
ListA
,
class
ListB
,
int
InCounter
,
int
OutCounter
>
struct
cl_arg_info_list_impl
;
template
<
class
Arg
,
class
...
Remaining
,
int
InCounter
,
int
OutCounter
>
struct
cl_arg_info_list_impl
<
detail
::
type_list
<>
,
detail
::
type_list
<
Arg
,
Remaining
...
>
,
InCounter
,
OutCounter
>
{
using
in_idx
=
in_index_of
<
InCounter
,
Arg
>
;
using
out_idx
=
out_index_of
<
OutCounter
,
Arg
>
;
using
type
=
typename
cl_arg_info_list_impl
<
detail
::
type_list
<
cl_arg_info
<
in_idx
::
value
,
out_idx
::
value
,
Arg
>>
,
detail
::
type_list
<
Remaining
...
>
,
in_idx
::
next
,
out_idx
::
next
>::
type
;
};
template
<
class
...
Args
,
class
Arg
,
class
...
Remaining
,
int
InCounter
,
int
OutCounter
>
struct
cl_arg_info_list_impl
<
detail
::
type_list
<
Args
...
>
,
detail
::
type_list
<
Arg
,
Remaining
...
>
,
InCounter
,
OutCounter
>
{
using
in_idx
=
in_index_of
<
InCounter
,
Arg
>
;
using
out_idx
=
out_index_of
<
OutCounter
,
Arg
>
;
using
type
=
typename
cl_arg_info_list_impl
<
detail
::
type_list
<
Args
...,
cl_arg_info
<
in_idx
::
value
,
out_idx
::
value
,
Arg
>>
,
detail
::
type_list
<
Remaining
...
>
,
in_idx
::
next
,
out_idx
::
next
>::
type
;
};
template
<
class
...
Args
,
int
InCounter
,
int
OutCounter
>
struct
cl_arg_info_list_impl
<
detail
::
type_list
<
Args
...
>
,
detail
::
type_list
<>
,
InCounter
,
OutCounter
>
{
using
type
=
detail
::
type_list
<
Args
...
>
;
};
template
<
class
List
>
struct
cl_arg_info_list
{
using
type
=
typename
cl_arg_info_list_impl
<
detail
::
type_list
<>
,
List
,
0
,
0
>::
type
;
};
/// Helpers for conversion in deprecated spawn functions
template
<
class
T
>
struct
to_input_arg
{
using
type
=
in
<
T
>
;
};
template
<
class
T
>
struct
to_output_arg
{
using
type
=
out
<
T
>
;
};
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/command.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <tuple>
#include <vector>
#include <numeric>
#include <algorithm>
#include <functional>
#include "caf/abstract_actor.hpp"
#include "caf/actor_cast.hpp"
#include "caf/logger.hpp"
#include "caf/raise_error.hpp"
#include "caf/response_promise.hpp"
#include "caf/detail/raw_ptr.hpp"
#include "caf/detail/scope_guard.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/nd_range.hpp"
#include "caf/opencl/arguments.hpp"
#include "caf/opencl/opencl_err.hpp"
namespace
caf
{
namespace
opencl
{
/// A command represents the execution of a kernel on a device. It handles the
/// OpenCL calls to enqueue the kernel with the index space and keeps references
/// to the management data during the execution. Furthermore, the command sends
/// the execution results to the responsible actor.
template
<
class
Actor
,
class
...
Ts
>
class
command
:
public
ref_counted
{
public:
using
result_types
=
detail
::
type_list
<
Ts
...
>
;
command
(
response_promise
promise
,
strong_actor_ptr
parent
,
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
detail
::
raw_mem_ptr
>
inputs
,
std
::
vector
<
detail
::
raw_mem_ptr
>
outputs
,
std
::
vector
<
detail
::
raw_mem_ptr
>
scratches
,
std
::
vector
<
size_t
>
lengths
,
message
msg
,
std
::
tuple
<
Ts
...
>
output_tuple
,
nd_range
range
)
:
lengths_
(
std
::
move
(
lengths
)),
promise_
(
std
::
move
(
promise
)),
cl_actor_
(
std
::
move
(
parent
)),
mem_in_events_
(
std
::
move
(
events
)),
input_buffers_
(
std
::
move
(
inputs
)),
output_buffers_
(
std
::
move
(
outputs
)),
scratch_buffers_
(
std
::
move
(
scratches
)),
results_
(
std
::
move
(
output_tuple
)),
msg_
(
std
::
move
(
msg
)),
range_
(
std
::
move
(
range
))
{
// nop
}
~
command
()
override
{
for
(
auto
&
e
:
mem_in_events_
)
{
if
(
e
)
v1callcl
(
CAF_CLF
(
clReleaseEvent
),
e
);
}
for
(
auto
&
e
:
mem_out_events_
)
{
if
(
e
)
v1callcl
(
CAF_CLF
(
clReleaseEvent
),
e
);
}
}
/// Enqueue the kernel for execution, schedule reading of the results and
/// set a callback to send the results to the actor identified by the handle.
/// Only called if the results includes at least one type that is not a
/// mem_ref.
template
<
class
Q
=
result_types
>
detail
::
enable_if_t
<!
detail
::
tl_forall
<
Q
,
is_ref_type
>::
value
>
enqueue
()
{
// Errors in this function can not be handled by opencl_err.hpp
// because they require non-standard error handling
CAF_LOG_TRACE
(
""
);
this
->
ref
();
// reference held by the OpenCL comand queue
auto
data_or_nullptr
=
[](
const
dim_vec
&
vec
)
{
return
vec
.
empty
()
?
nullptr
:
vec
.
data
();
};
auto
parent
=
static_cast
<
Actor
*>
(
actor_cast
<
abstract_actor
*>
(
cl_actor_
));
// OpenCL expects cl_uint (unsigned int), hence the cast
mem_out_events_
.
emplace_back
();
auto
success
=
invoke_cl
(
clEnqueueNDRangeKernel
,
parent
->
queue_
.
get
(),
parent
->
kernel_
.
get
(),
static_cast
<
unsigned
int
>
(
range_
.
dimensions
().
size
()),
data_or_nullptr
(
range_
.
offsets
()),
data_or_nullptr
(
range_
.
dimensions
()),
data_or_nullptr
(
range_
.
local_dimensions
()),
static_cast
<
unsigned
int
>
(
mem_in_events_
.
size
()),
(
mem_in_events_
.
empty
()
?
nullptr
:
mem_in_events_
.
data
()),
&
mem_out_events_
.
back
()
);
if
(
!
success
)
return
;
size_t
pos
=
0
;
CAF_ASSERT
(
!
mem_out_events_
.
empty
());
enqueue_read_buffers
(
pos
,
mem_out_events_
,
detail
::
get_indices
(
results_
));
cl_event
marker_event
;
#if defined(__APPLE__)
success
=
invoke_cl
(
clEnqueueMarkerWithWaitList
,
parent
->
queue_
.
get
(),
static_cast
<
unsigned
int
>
(
mem_out_events_
.
size
()),
mem_out_events_
.
data
(),
&
marker_event
);
#else
success
=
invoke_cl
(
clEnqueueMarker
,
parent
->
queue_
.
get
(),
&
marker_event
);
#endif
callback_
.
reset
(
marker_event
,
false
);
if
(
!
success
)
return
;
auto
cb
=
[](
cl_event
,
cl_int
,
void
*
data
)
{
auto
cmd
=
reinterpret_cast
<
command
*>
(
data
);
cmd
->
handle_results
();
cmd
->
deref
();
};
if
(
!
invoke_cl
(
clSetEventCallback
,
callback_
.
get
(),
CL_COMPLETE
,
std
::
move
(
cb
),
this
))
return
;
v3callcl
(
clFlush
,
parent
->
queue_
.
get
());
}
/// Enqueue the kernel for execution and send the mem_refs relating to the
/// results to the next actor. A callback is set to clean up the commmand
/// once the execution is finished. Only called if the results only consist
/// of mem_ref types.
template
<
class
Q
=
result_types
>
detail
::
enable_if_t
<
detail
::
tl_forall
<
Q
,
is_ref_type
>::
value
>
enqueue
()
{
// Errors in this function can not be handled by opencl_err.hpp
// because they require non-standard error handling
CAF_LOG_TRACE
(
""
);
this
->
ref
();
// reference held by the OpenCL command queue
auto
data_or_nullptr
=
[](
const
dim_vec
&
vec
)
{
return
vec
.
empty
()
?
nullptr
:
vec
.
data
();
};
auto
parent
=
static_cast
<
Actor
*>
(
actor_cast
<
abstract_actor
*>
(
cl_actor_
));
cl_event
execution_event
;
auto
success
=
invoke_cl
(
clEnqueueNDRangeKernel
,
parent
->
queue_
.
get
(),
parent
->
kernel_
.
get
(),
static_cast
<
cl_uint
>
(
range_
.
dimensions
().
size
()),
data_or_nullptr
(
range_
.
offsets
()),
data_or_nullptr
(
range_
.
dimensions
()),
data_or_nullptr
(
range_
.
local_dimensions
()),
static_cast
<
unsigned
int
>
(
mem_in_events_
.
size
()),
(
mem_in_events_
.
empty
()
?
nullptr
:
mem_in_events_
.
data
()),
&
execution_event
);
callback_
.
reset
(
execution_event
,
false
);
if
(
!
success
)
return
;
auto
cb
=
[](
cl_event
,
cl_int
,
void
*
data
)
{
auto
c
=
reinterpret_cast
<
command
*>
(
data
);
c
->
deref
();
};
if
(
!
invoke_cl
(
clSetEventCallback
,
callback_
.
get
(),
CL_COMPLETE
,
std
::
move
(
cb
),
this
))
return
;
v3callcl
(
clFlush
,
parent
->
queue_
.
get
());
auto
msg
=
msg_adding_event
{
callback_
}(
results_
);
promise_
.
deliver
(
std
::
move
(
msg
));
}
private:
template
<
long
I
,
class
T
>
void
enqueue_read
(
std
::
vector
<
T
>&
,
std
::
vector
<
cl_event
>&
events
,
size_t
&
pos
)
{
auto
p
=
static_cast
<
Actor
*>
(
actor_cast
<
abstract_actor
*>
(
cl_actor_
));
events
.
emplace_back
();
auto
size
=
lengths_
[
pos
];
auto
buffer_size
=
sizeof
(
T
)
*
size
;
std
::
get
<
I
>
(
results_
).
resize
(
size
);
auto
err
=
clEnqueueReadBuffer
(
p
->
queue_
.
get
(),
output_buffers_
[
pos
].
get
(),
CL_FALSE
,
0
,
buffer_size
,
std
::
get
<
I
>
(
results_
).
data
(),
1
,
events
.
data
(),
&
events
.
back
());
if
(
err
!=
CL_SUCCESS
)
{
this
->
deref
();
// failed to enqueue command
CAF_RAISE_ERROR
(
"failed to enqueue command"
);
}
pos
+=
1
;
}
template
<
long
I
,
class
T
>
void
enqueue_read
(
mem_ref
<
T
>&
,
std
::
vector
<
cl_event
>&
,
size_t
&
)
{
// Nothing to read back if we return references.
}
void
enqueue_read_buffers
(
size_t
&
,
std
::
vector
<
cl_event
>&
,
detail
::
int_list
<>
)
{
// end of recursion
}
template
<
long
I
,
long
...
Is
>
void
enqueue_read_buffers
(
size_t
&
pos
,
std
::
vector
<
cl_event
>&
events
,
detail
::
int_list
<
I
,
Is
...
>
)
{
enqueue_read
<
I
>
(
std
::
get
<
I
>
(
results_
),
events
,
pos
);
enqueue_read_buffers
(
pos
,
events
,
detail
::
int_list
<
Is
...
>
{});
}
// handle results if execution result includes a value type
void
handle_results
()
{
auto
parent
=
static_cast
<
Actor
*>
(
actor_cast
<
abstract_actor
*>
(
cl_actor_
));
auto
&
map_fun
=
parent
->
map_results_
;
auto
msg
=
map_fun
?
apply_args
(
map_fun
,
detail
::
get_indices
(
results_
),
results_
)
:
message_from_results
{}(
results_
);
promise_
.
deliver
(
std
::
move
(
msg
));
}
// call function F and derefenrence the command on failure
template
<
class
F
,
class
...
Us
>
bool
invoke_cl
(
F
f
,
Us
&&
...
xs
)
{
auto
err
=
f
(
std
::
forward
<
Us
>
(
xs
)...);
if
(
err
==
CL_SUCCESS
)
return
true
;
CAF_LOG_ERROR
(
"error: "
<<
opencl_error
(
err
));
this
->
deref
();
return
false
;
}
std
::
vector
<
size_t
>
lengths_
;
response_promise
promise_
;
strong_actor_ptr
cl_actor_
;
std
::
vector
<
cl_event
>
mem_in_events_
;
std
::
vector
<
cl_event
>
mem_out_events_
;
detail
::
raw_event_ptr
callback_
;
std
::
vector
<
detail
::
raw_mem_ptr
>
input_buffers_
;
std
::
vector
<
detail
::
raw_mem_ptr
>
output_buffers_
;
std
::
vector
<
detail
::
raw_mem_ptr
>
scratch_buffers_
;
std
::
tuple
<
Ts
...
>
results_
;
message
msg_
;
// keeps the argument buffers alive for async copy to device
nd_range
range_
;
};
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/device.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <vector>
#include "caf/sec.hpp"
#include "caf/detail/raw_ptr.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/opencl_err.hpp"
namespace
caf
{
namespace
opencl
{
class
program
;
class
manager
;
template
<
class
T
>
class
mem_ref
;
class
device
;
using
device_ptr
=
intrusive_ptr
<
device
>
;
class
device
:
public
ref_counted
{
public:
friend
class
program
;
friend
class
manager
;
template
<
class
T
>
friend
class
mem_ref
;
template
<
class
T
,
class
...
Ts
>
friend
intrusive_ptr
<
T
>
caf
::
make_counted
(
Ts
&&
...);
~
device
();
/// Create an argument for an OpenCL kernel with data placed in global memory.
template
<
class
T
>
mem_ref
<
T
>
global_argument
(
const
std
::
vector
<
T
>&
data
,
cl_mem_flags
flags
=
buffer_type
::
input_output
,
optional
<
size_t
>
size
=
none
,
cl_bool
blocking
=
CL_FALSE
)
{
size_t
num_elements
=
size
?
*
size
:
data
.
size
();
size_t
buffer_size
=
sizeof
(
T
)
*
num_elements
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
flags
,
buffer_size
,
nullptr
);
detail
::
raw_event_ptr
event
{
v1get
<
cl_event
>
(
CAF_CLF
(
clEnqueueWriteBuffer
),
queue_
.
get
(),
buffer
,
blocking
,
cl_uint
{
0
},
buffer_size
,
data
.
data
()),
false
};
return
mem_ref
<
T
>
{
num_elements
,
queue_
,
std
::
move
(
buffer
),
flags
,
std
::
move
(
event
)};
}
/// Create an argument for an OpenCL kernel in global memory without data.
template
<
class
T
>
mem_ref
<
T
>
scratch_argument
(
size_t
size
,
cl_mem_flags
flags
=
buffer_type
::
scratch_space
)
{
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
flags
,
sizeof
(
T
)
*
size
,
nullptr
);
return
mem_ref
<
T
>
{
size
,
queue_
,
std
::
move
(
buffer
),
flags
,
nullptr
};
}
template
<
class
T
>
expected
<
mem_ref
<
T
>>
copy
(
mem_ref
<
T
>&
mem
)
{
if
(
!
mem
.
get
())
return
make_error
(
sec
::
runtime_error
,
"No memory assigned."
);
auto
buffer_size
=
sizeof
(
T
)
*
mem
.
size
();
cl_event
event
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
mem
.
access
(),
buffer_size
,
nullptr
);
std
::
vector
<
cl_event
>
prev_events
;
cl_event
e
=
mem
.
take_event
();
if
(
e
)
prev_events
.
push_back
(
e
);
auto
err
=
clEnqueueCopyBuffer
(
queue_
.
get
(),
mem
.
get
().
get
(),
buffer
,
0
,
0
,
// no offset for now
buffer_size
,
prev_events
.
size
(),
prev_events
.
data
(),
&
event
);
if
(
err
!=
CL_SUCCESS
)
return
make_error
(
sec
::
runtime_error
,
opencl_error
(
err
));
// callback to release the previous event
if
(
e
)
{
err
=
clSetEventCallback
(
event
,
CL_COMPLETE
,
[](
cl_event
,
cl_int
,
void
*
data
)
{
auto
tmp
=
reinterpret_cast
<
cl_event
>
(
data
);
if
(
tmp
)
clReleaseEvent
(
tmp
);
},
e
);
if
(
err
!=
CL_SUCCESS
)
return
make_error
(
sec
::
runtime_error
,
opencl_error
(
err
));
}
// decrements the previous event we used for waiting above
return
mem_ref
<
T
>
(
mem
.
size
(),
queue_
,
std
::
move
(
buffer
),
mem
.
access
(),
{
event
,
false
});
}
/// Initialize a new device in a context using a specific device_id
static
device_ptr
create
(
const
detail
::
raw_context_ptr
&
context
,
const
detail
::
raw_device_ptr
&
device_id
,
unsigned
id
);
/// Synchronizes all commands in its queue, waiting for them to finish.
void
synchronize
();
/// Get the id assigned by caf
inline
unsigned
id
()
const
;
/// Returns device info on CL_DEVICE_ADDRESS_BITS
inline
cl_uint
address_bits
()
const
;
/// Returns device info on CL_DEVICE_ENDIAN_LITTLE
inline
cl_bool
little_endian
()
const
;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
inline
cl_ulong
global_mem_cache_size
()
const
;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
inline
cl_uint
global_mem_cacheline_size
()
const
;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_SIZE
inline
cl_ulong
global_mem_size
()
const
;
/// Returns device info on CL_DEVICE_HOST_UNIFIED_MEMORY
inline
cl_bool
host_unified_memory
()
const
;
/// Returns device info on CL_DEVICE_LOCAL_MEM_SIZE
inline
cl_ulong
local_mem_size
()
const
;
/// Returns device info on CL_DEVICE_LOCAL_MEM_TYPE
inline
cl_ulong
local_mem_type
()
const
;
/// Returns device info on CL_DEVICE_MAX_CLOCK_FREQUENCY
inline
cl_uint
max_clock_frequency
()
const
;
/// Returns device info on CL_DEVICE_MAX_COMPUTE_UNITS
inline
cl_uint
max_compute_units
()
const
;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_ARGS
inline
cl_uint
max_constant_args
()
const
;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
inline
cl_ulong
max_constant_buffer_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_MEM_ALLOC_SIZE
inline
cl_ulong
max_mem_alloc_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_PARAMETER_SIZE
inline
size_t
max_parameter_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_WORK_GROUP_SIZE
inline
size_t
max_work_group_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
inline
cl_uint
max_work_item_dimensions
()
const
;
/// Returns device info on CL_DEVICE_PROFILING_TIMER_RESOLUTION
inline
size_t
profiling_timer_resolution
()
const
;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_SIZES
inline
const
dim_vec
&
max_work_item_sizes
()
const
;
/// Returns device info on CL_DEVICE_TYPE
inline
device_type
type
()
const
;
/// Returns device info on CL_DEVICE_EXTENSIONS
inline
const
std
::
vector
<
std
::
string
>&
extensions
()
const
;
/// Returns device info on CL_DEVICE_OPENCL_C_VERSION
inline
const
std
::
string
&
opencl_c_version
()
const
;
/// Returns device info on CL_DEVICE_VENDOR
inline
const
std
::
string
&
device_vendor
()
const
;
/// Returns device info on CL_DEVICE_VERSION
inline
const
std
::
string
&
device_version
()
const
;
/// Returns device info on CL_DRIVER_VERSION
inline
const
std
::
string
&
driver_version
()
const
;
/// Returns device info on CL_DEVICE_NAME
inline
const
std
::
string
&
name
()
const
;
private:
device
(
detail
::
raw_device_ptr
device_id
,
detail
::
raw_command_queue_ptr
queue
,
detail
::
raw_context_ptr
context
,
unsigned
id
);
template
<
class
T
>
static
T
info
(
const
detail
::
raw_device_ptr
&
device_id
,
unsigned
info_flag
)
{
T
value
;
clGetDeviceInfo
(
device_id
.
get
(),
info_flag
,
sizeof
(
T
),
&
value
,
nullptr
);
return
value
;
}
static
std
::
string
info_string
(
const
detail
::
raw_device_ptr
&
device_id
,
unsigned
info_flag
);
detail
::
raw_device_ptr
device_id_
;
detail
::
raw_command_queue_ptr
queue_
;
detail
::
raw_context_ptr
context_
;
unsigned
id_
;
bool
profiling_enabled_
;
// CL_DEVICE_QUEUE_PROPERTIES
bool
out_of_order_execution_
;
// CL_DEVICE_QUEUE_PROPERTIES
cl_uint
address_bits_
;
// CL_DEVICE_ADDRESS_BITS
cl_bool
little_endian_
;
// CL_DEVICE_ENDIAN_LITTLE
cl_ulong
global_mem_cache_size_
;
// CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
cl_uint
global_mem_cacheline_size_
;
// CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
cl_ulong
global_mem_size_
;
// CL_DEVICE_GLOBAL_MEM_SIZE
cl_bool
host_unified_memory_
;
// CL_DEVICE_HOST_UNIFIED_MEMORY
cl_ulong
local_mem_size_
;
// CL_DEVICE_LOCAL_MEM_SIZE
cl_uint
local_mem_type_
;
// CL_DEVICE_LOCAL_MEM_TYPE
cl_uint
max_clock_frequency_
;
// CL_DEVICE_MAX_CLOCK_FREQUENCY
cl_uint
max_compute_units_
;
// CL_DEVICE_MAX_COMPUTE_UNITS
cl_uint
max_constant_args_
;
// CL_DEVICE_MAX_CONSTANT_ARGS
cl_ulong
max_constant_buffer_size_
;
// CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
cl_ulong
max_mem_alloc_size_
;
// CL_DEVICE_MAX_MEM_ALLOC_SIZE
size_t
max_parameter_size_
;
// CL_DEVICE_MAX_PARAMETER_SIZE
size_t
max_work_group_size_
;
// CL_DEVICE_MAX_WORK_GROUP_SIZE
cl_uint
max_work_item_dimensions_
;
// CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
size_t
profiling_timer_resolution_
;
// CL_DEVICE_PROFILING_TIMER_RESOLUTION
dim_vec
max_work_item_sizes_
;
// CL_DEVICE_MAX_WORK_ITEM_SIZES
device_type
device_type_
;
// CL_DEVICE_TYPE
std
::
vector
<
std
::
string
>
extensions_
;
// CL_DEVICE_EXTENSIONS
std
::
string
opencl_c_version_
;
// CL_DEVICE_OPENCL_C_VERSION
std
::
string
device_vendor_
;
// CL_DEVICE_VENDOR
std
::
string
device_version_
;
// CL_DEVICE_VERSION
std
::
string
driver_version_
;
// CL_DRIVER_VERSION
std
::
string
name_
;
// CL_DEVICE_NAME
};
/******************************************************************************\
* implementation of inline member functions *
\******************************************************************************/
inline
unsigned
device
::
id
()
const
{
return
id_
;
}
inline
cl_uint
device
::
address_bits
()
const
{
return
address_bits_
;
}
inline
cl_bool
device
::
little_endian
()
const
{
return
little_endian_
;
}
inline
cl_ulong
device
::
global_mem_cache_size
()
const
{
return
global_mem_cache_size_
;
}
inline
cl_uint
device
::
global_mem_cacheline_size
()
const
{
return
global_mem_cacheline_size_
;
}
inline
cl_ulong
device
::
global_mem_size
()
const
{
return
global_mem_size_
;
}
inline
cl_bool
device
::
host_unified_memory
()
const
{
return
host_unified_memory_
;
}
inline
cl_ulong
device
::
local_mem_size
()
const
{
return
local_mem_size_
;
}
inline
cl_ulong
device
::
local_mem_type
()
const
{
return
local_mem_size_
;
}
inline
cl_uint
device
::
max_clock_frequency
()
const
{
return
max_clock_frequency_
;
}
inline
cl_uint
device
::
max_compute_units
()
const
{
return
max_compute_units_
;
}
inline
cl_uint
device
::
max_constant_args
()
const
{
return
max_constant_args_
;
}
inline
cl_ulong
device
::
max_constant_buffer_size
()
const
{
return
max_constant_buffer_size_
;
}
inline
cl_ulong
device
::
max_mem_alloc_size
()
const
{
return
max_mem_alloc_size_
;
}
inline
size_t
device
::
max_parameter_size
()
const
{
return
max_parameter_size_
;
}
inline
size_t
device
::
max_work_group_size
()
const
{
return
max_work_group_size_
;
}
inline
cl_uint
device
::
max_work_item_dimensions
()
const
{
return
max_work_item_dimensions_
;
}
inline
size_t
device
::
profiling_timer_resolution
()
const
{
return
profiling_timer_resolution_
;
}
inline
const
dim_vec
&
device
::
max_work_item_sizes
()
const
{
return
max_work_item_sizes_
;
}
inline
device_type
device
::
type
()
const
{
return
device_type_
;
}
inline
const
std
::
vector
<
std
::
string
>&
device
::
extensions
()
const
{
return
extensions_
;
}
inline
const
std
::
string
&
device
::
opencl_c_version
()
const
{
return
opencl_c_version_
;
}
inline
const
std
::
string
&
device
::
device_vendor
()
const
{
return
device_vendor_
;
}
inline
const
std
::
string
&
device
::
device_version
()
const
{
return
device_version_
;
}
inline
const
std
::
string
&
device
::
driver_version
()
const
{
return
driver_version_
;
}
inline
const
std
::
string
&
device
::
name
()
const
{
return
name_
;
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/global.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <string>
#include <iostream>
#include "caf/config.hpp"
#include "caf/detail/limited_vector.hpp"
#if defined(CAF_MACOS) || defined(CAF_IOS)
# include <OpenCL/opencl.h>
#else
# include <CL/opencl.h>
#endif
// needed for OpenCL 1.0 compatibility (works around missing clReleaseDevice)
extern
"C"
{
cl_int
clReleaseDeviceDummy
(
cl_device_id
);
cl_int
clRetainDeviceDummy
(
cl_device_id
);
}
// extern "C"
namespace
caf
{
namespace
opencl
{
enum
device_type
{
def
=
CL_DEVICE_TYPE_DEFAULT
,
cpu
=
CL_DEVICE_TYPE_CPU
,
gpu
=
CL_DEVICE_TYPE_GPU
,
accelerator
=
CL_DEVICE_TYPE_ACCELERATOR
,
custom
=
CL_DEVICE_TYPE_CUSTOM
,
all
=
CL_DEVICE_TYPE_ALL
};
/// Default values to create OpenCL buffers
enum
buffer_type
:
cl_mem_flags
{
input
=
CL_MEM_READ_WRITE
|
CL_MEM_HOST_WRITE_ONLY
,
input_output
=
CL_MEM_READ_WRITE
,
output
=
CL_MEM_READ_WRITE
|
CL_MEM_HOST_READ_ONLY
,
scratch_space
=
CL_MEM_READ_WRITE
|
CL_MEM_HOST_NO_ACCESS
};
std
::
ostream
&
operator
<<
(
std
::
ostream
&
os
,
device_type
dev
);
device_type
device_type_from_ulong
(
cl_ulong
dev
);
/// A vector of up to three elements used for OpenCL dimensions.
using
dim_vec
=
detail
::
limited_vector
<
size_t
,
3
>
;
std
::
string
opencl_error
(
cl_int
err
);
std
::
string
event_status
(
cl_event
event
);
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/manager.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <atomic>
#include <vector>
#include <algorithm>
#include <functional>
#include "caf/optional.hpp"
#include "caf/config.hpp"
#include "caf/actor_system.hpp"
#include "caf/detail/raw_ptr.hpp"
#include "caf/detail/spawn_helper.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/actor_facade.hpp"
namespace
caf
{
namespace
opencl
{
class
manager
:
public
actor_system
::
module
{
public:
friend
class
program
;
friend
class
actor_system
;
friend
detail
::
raw_command_queue_ptr
command_queue
(
uint32_t
id
);
manager
(
const
manager
&
)
=
delete
;
manager
&
operator
=
(
const
manager
&
)
=
delete
;
/// Get the device with id, which is assigned sequientally.
optional
<
device_ptr
>
find_device
(
size_t
dev_id
=
0
)
const
;
/// Get the first device that satisfies the predicate.
/// The predicate should accept a `const device&` and return a bool;
template
<
class
UnaryPredicate
>
optional
<
device_ptr
>
find_device_if
(
UnaryPredicate
p
)
const
{
for
(
auto
&
pl
:
platforms_
)
{
for
(
auto
&
dev
:
pl
->
devices
())
{
if
(
p
(
dev
))
return
dev
;
}
}
return
none
;
}
void
start
()
override
;
void
stop
()
override
;
void
init
(
actor_system_config
&
)
override
;
id_t
id
()
const
override
;
void
*
subtype_ptr
()
override
;
static
actor_system
::
module
*
make
(
actor_system
&
sys
,
detail
::
type_list
<>
);
// OpenCL functionality
/// @brief Factory method, that creates a caf::opencl::program
/// reading the source from given @p path.
/// @returns A program object.
program_ptr
create_program_from_file
(
const
char
*
path
,
const
char
*
options
=
nullptr
,
uint32_t
device_id
=
0
);
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
program_ptr
create_program
(
const
char
*
kernel_source
,
const
char
*
options
=
nullptr
,
uint32_t
device_id
=
0
);
/// @brief Factory method, that creates a caf::opencl::program
/// reading the source from given @p path.
/// @returns A program object.
program_ptr
create_program_from_file
(
const
char
*
path
,
const
char
*
options
,
const
device_ptr
dev
);
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
program_ptr
create_program
(
const
char
*
kernel_source
,
const
char
*
options
,
const
device_ptr
dev
);
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template
<
class
T
,
class
...
Ts
>
detail
::
enable_if_t
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>
spawn
(
const
opencl
::
program_ptr
prog
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
T
&&
x
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
false
,
T
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
prog
,
fname
,
range
,
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template
<
class
T
,
class
...
Ts
>
detail
::
enable_if_t
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>
spawn
(
const
char
*
source
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
T
&&
x
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
false
,
T
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
range
,
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template
<
class
Fun
,
class
...
Ts
>
actor
spawn
(
const
opencl
::
program_ptr
prog
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
message
&
)
>
map_args
,
Fun
map_result
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
false
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
prog
,
fname
,
range
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template
<
class
Fun
,
class
...
Ts
>
actor
spawn
(
const
char
*
source
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
message
&
)
>
map_args
,
Fun
map_result
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
false
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
range
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward
<
Ts
>
(
xs
)...);
}
// --- Only accept the input mapping function ---
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template
<
class
T
,
class
...
Ts
>
detail
::
enable_if_t
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>
spawn
(
const
opencl
::
program_ptr
prog
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
message
&
)
>
map_args
,
T
&&
x
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
false
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
prog
,
fname
,
range
,
std
::
move
(
map_args
),
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template
<
class
T
,
class
...
Ts
>
detail
::
enable_if_t
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>
spawn
(
const
char
*
source
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
message
&
)
>
map_args
,
T
&&
x
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
false
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
range
,
std
::
move
(
map_args
),
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
// --- Input mapping function accepts config as well ---
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template
<
class
Fun
,
class
...
Ts
>
detail
::
enable_if_t
<!
opencl
::
is_opencl_arg
<
Fun
>::
value
,
actor
>
spawn
(
const
opencl
::
program_ptr
prog
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
nd_range
&
,
message
&
)
>
map_args
,
Fun
map_result
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
true
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
prog
,
fname
,
range
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template
<
class
Fun
,
class
...
Ts
>
detail
::
enable_if_t
<!
opencl
::
is_opencl_arg
<
Fun
>::
value
,
actor
>
spawn
(
const
char
*
source
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
nd_range
&
,
message
&
)
>
map_args
,
Fun
map_result
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
true
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
range
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template
<
class
T
,
class
...
Ts
>
detail
::
enable_if_t
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>
spawn
(
const
opencl
::
program_ptr
prog
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
nd_range
&
,
message
&
)
>
map_args
,
T
&&
x
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
true
,
T
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
prog
,
fname
,
range
,
std
::
move
(
map_args
),
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template
<
class
T
,
class
...
Ts
>
detail
::
enable_if_t
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>
spawn
(
const
char
*
source
,
const
char
*
fname
,
const
opencl
::
nd_range
&
range
,
std
::
function
<
optional
<
message
>
(
nd_range
&
,
message
&
)
>
map_args
,
T
&&
x
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
true
,
T
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
range
,
std
::
move
(
map_args
),
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
protected:
manager
(
actor_system
&
sys
);
~
manager
()
override
;
private:
actor_system
&
system_
;
std
::
vector
<
platform_ptr
>
platforms_
;
};
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/mem_ref.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <ios>
#include <vector>
#include "caf/sec.hpp"
#include "caf/optional.hpp"
#include "caf/ref_counted.hpp"
#include "caf/detail/raw_ptr.hpp"
namespace
caf
{
namespace
opencl
{
// Tag to separate mem_refs from other types in messages.
struct
ref_tag
{};
class
device
;
/// A reference type for buffers on a OpenCL devive. Access is not thread safe.
/// Hence, a mem_ref should only be passed to actors sequentially.
template
<
class
T
>
class
mem_ref
:
ref_tag
{
public:
using
value_type
=
T
;
friend
struct
msg_adding_event
;
template
<
bool
PassConfig
,
class
...
Ts
>
friend
class
actor_facade
;
friend
class
device
;
expected
<
std
::
vector
<
T
>>
data
(
optional
<
size_t
>
result_size
=
none
)
{
if
(
!
memory_
)
return
make_error
(
sec
::
runtime_error
,
"No memory assigned."
);
if
(
0
!=
(
access_
&
CL_MEM_HOST_NO_ACCESS
))
return
make_error
(
sec
::
runtime_error
,
"No memory access."
);
if
(
result_size
&&
*
result_size
>
num_elements_
)
return
make_error
(
sec
::
runtime_error
,
"Buffer has less elements."
);
auto
num_elements
=
(
result_size
?
*
result_size
:
num_elements_
);
auto
buffer_size
=
sizeof
(
T
)
*
num_elements
;
std
::
vector
<
T
>
buffer
(
num_elements
);
std
::
vector
<
cl_event
>
prev_events
;
if
(
event_
)
prev_events
.
push_back
(
event_
.
get
());
cl_event
event
;
auto
err
=
clEnqueueReadBuffer
(
queue_
.
get
(),
memory_
.
get
(),
CL_TRUE
,
0
,
buffer_size
,
buffer
.
data
(),
static_cast
<
cl_uint
>
(
prev_events
.
size
()),
prev_events
.
data
(),
&
event
);
if
(
err
!=
CL_SUCCESS
)
return
make_error
(
sec
::
runtime_error
,
opencl_error
(
err
));
// decrements the previous event we used for waiting above
event_
.
reset
(
event
,
false
);
return
buffer
;
}
void
reset
()
{
num_elements_
=
0
;
access_
=
CL_MEM_HOST_NO_ACCESS
;
memory_
.
reset
();
access_
=
0
;
event_
.
reset
();
}
inline
const
detail
::
raw_mem_ptr
&
get
()
const
{
return
memory_
;
}
inline
size_t
size
()
const
{
return
num_elements_
;
}
inline
cl_mem_flags
access
()
const
{
return
access_
;
}
mem_ref
()
:
num_elements_
{
0
},
access_
{
CL_MEM_HOST_NO_ACCESS
},
queue_
{
nullptr
},
event_
{
nullptr
},
memory_
{
nullptr
}
{
// nop
}
mem_ref
(
size_t
num_elements
,
detail
::
raw_command_queue_ptr
queue
,
detail
::
raw_mem_ptr
memory
,
cl_mem_flags
access
,
detail
::
raw_event_ptr
event
)
:
num_elements_
{
num_elements
},
access_
{
access
},
queue_
{
queue
},
event_
{
event
},
memory_
{
memory
}
{
// nop
}
mem_ref
(
size_t
num_elements
,
detail
::
raw_command_queue_ptr
queue
,
cl_mem
memory
,
cl_mem_flags
access
,
detail
::
raw_event_ptr
event
)
:
num_elements_
{
num_elements
},
access_
{
access
},
queue_
{
queue
},
event_
{
event
},
memory_
{
memory
}
{
// nop
}
mem_ref
(
mem_ref
&&
other
)
=
default
;
mem_ref
(
const
mem_ref
&
other
)
=
default
;
mem_ref
&
operator
=
(
mem_ref
<
T
>&&
other
)
=
default
;
mem_ref
&
operator
=
(
const
mem_ref
&
other
)
=
default
;
~
mem_ref
()
{
// nop
}
private:
inline
void
set_event
(
cl_event
e
,
bool
increment_reference
=
true
)
{
event_
.
reset
(
e
,
increment_reference
);
}
inline
void
set_event
(
detail
::
raw_event_ptr
e
)
{
event_
=
std
::
move
(
e
);
}
inline
detail
::
raw_event_ptr
event
()
{
return
event_
;
}
inline
cl_event
take_event
()
{
return
event_
.
release
();
}
size_t
num_elements_
;
cl_mem_flags
access_
;
detail
::
raw_command_queue_ptr
queue_
;
detail
::
raw_event_ptr
event_
;
detail
::
raw_mem_ptr
memory_
;
};
/// Updates the reference types in a message with a given event.
struct
msg_adding_event
{
msg_adding_event
(
detail
::
raw_event_ptr
event
)
:
event_
(
event
)
{
// nop
}
template
<
class
T
,
class
...
Ts
>
message
operator
()(
T
&
x
,
Ts
&
...
xs
)
{
return
make_message
(
add_event
(
std
::
move
(
x
)),
add_event
(
std
::
move
(
xs
))...);
}
template
<
class
...
Ts
>
message
operator
()(
std
::
tuple
<
Ts
...
>&
values
)
{
return
apply_args
(
*
this
,
detail
::
get_indices
(
values
),
values
);
}
template
<
class
T
>
mem_ref
<
T
>
add_event
(
mem_ref
<
T
>
ref
)
{
ref
.
set_event
(
event_
);
return
std
::
move
(
ref
);
}
detail
::
raw_event_ptr
event_
;
};
}
// namespace opencl
template
<
class
T
>
struct
allowed_unsafe_message_type
<
opencl
::
mem_ref
<
T
>>
:
std
::
true_type
{};
}
// namespace caf
libcaf_opencl/caf/opencl/nd_range.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include "caf/opencl/global.hpp"
namespace
caf
{
namespace
opencl
{
class
nd_range
{
public:
nd_range
(
const
opencl
::
dim_vec
&
dimensions
,
const
opencl
::
dim_vec
&
offsets
=
{},
const
opencl
::
dim_vec
&
local_dimensions
=
{})
:
dims_
{
dimensions
},
offset_
{
offsets
},
local_dims_
{
local_dimensions
}
{
// nop
}
nd_range
(
opencl
::
dim_vec
&&
dimensions
,
opencl
::
dim_vec
&&
offsets
=
{},
opencl
::
dim_vec
&&
local_dimensions
=
{})
:
dims_
{
std
::
move
(
dimensions
)},
offset_
{
std
::
move
(
offsets
)},
local_dims_
{
std
::
move
(
local_dimensions
)}
{
// nop
}
nd_range
(
const
nd_range
&
)
=
default
;
nd_range
(
nd_range
&&
)
=
default
;
nd_range
&
operator
=
(
const
nd_range
&
)
=
default
;
nd_range
&
operator
=
(
nd_range
&&
)
=
default
;
const
opencl
::
dim_vec
&
dimensions
()
const
{
return
dims_
;
}
const
opencl
::
dim_vec
&
offsets
()
const
{
return
offset_
;
}
const
opencl
::
dim_vec
&
local_dimensions
()
const
{
return
local_dims_
;
}
private:
opencl
::
dim_vec
dims_
;
opencl
::
dim_vec
offset_
;
opencl
::
dim_vec
local_dims_
;
};
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/opencl_err.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Raphael Hiesgen <raphael.hiesgen (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 LICENSE_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. *
******************************************************************************/
#pragma once
#include "caf/opencl/global.hpp"
#include "caf/logger.hpp"
#define CAF_CLF(funname) #funname , funname
namespace
caf
{
namespace
opencl
{
void
throwcl
(
const
char
*
fname
,
cl_int
err
);
void
CL_CALLBACK
pfn_notify
(
const
char
*
errinfo
,
const
void
*
,
size_t
,
void
*
);
// call convention for simply calling a function
template
<
class
F
,
class
...
Ts
>
void
v1callcl
(
const
char
*
fname
,
F
f
,
Ts
&&
...
vs
)
{
throwcl
(
fname
,
f
(
std
::
forward
<
Ts
>
(
vs
)...));
}
// call convention for simply calling a function, not using the last argument
template
<
class
F
,
class
...
Ts
>
void
v2callcl
(
const
char
*
fname
,
F
f
,
Ts
&&
...
vs
)
{
throwcl
(
fname
,
f
(
std
::
forward
<
Ts
>
(
vs
)...,
nullptr
));
}
// call convention for simply calling a function, and logging errors
template
<
class
F
,
class
...
Ts
>
void
v3callcl
(
F
f
,
Ts
&&
...
vs
)
{
auto
err
=
f
(
std
::
forward
<
Ts
>
(
vs
)...);
if
(
err
!=
CL_SUCCESS
)
CAF_LOG_ERROR
(
"error: "
<<
opencl_error
(
err
));
}
// call convention with `result` argument at the end returning `err`, not
// using the second last argument (set to nullptr) nor the one before (set to 0)
template
<
class
R
,
class
F
,
class
...
Ts
>
R
v1get
(
const
char
*
fname
,
F
f
,
Ts
&&
...
vs
)
{
R
result
;
throwcl
(
fname
,
f
(
std
::
forward
<
Ts
>
(
vs
)...,
cl_uint
{
0
},
nullptr
,
&
result
));
return
result
;
}
// call convention with `err` argument at the end returning `result`
template
<
class
F
,
class
...
Ts
>
auto
v2get
(
const
char
*
fname
,
F
f
,
Ts
&&
...
vs
)
->
decltype
(
f
(
std
::
forward
<
Ts
>
(
vs
)...,
nullptr
))
{
cl_int
err
;
auto
result
=
f
(
std
::
forward
<
Ts
>
(
vs
)...,
&
err
);
throwcl
(
fname
,
err
);
return
result
;
}
// call convention with `result` argument at second last position (preceeded by
// its size) followed by an ingored void* argument (nullptr) returning `err`
template
<
class
R
,
class
F
,
class
...
Ts
>
R
v3get
(
const
char
*
fname
,
F
f
,
Ts
&&
...
vs
)
{
R
result
;
throwcl
(
fname
,
f
(
std
::
forward
<
Ts
>
(
vs
)...,
sizeof
(
R
),
&
result
,
nullptr
));
return
result
;
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/platform.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include "caf/ref_counted.hpp"
#include "caf/opencl/device.hpp"
namespace
caf
{
namespace
opencl
{
class
platform
;
using
platform_ptr
=
intrusive_ptr
<
platform
>
;
class
platform
:
public
ref_counted
{
public:
friend
class
program
;
template
<
class
T
,
class
...
Ts
>
friend
intrusive_ptr
<
T
>
caf
::
make_counted
(
Ts
&&
...);
inline
const
std
::
vector
<
device_ptr
>&
devices
()
const
;
inline
const
std
::
string
&
name
()
const
;
inline
const
std
::
string
&
vendor
()
const
;
inline
const
std
::
string
&
version
()
const
;
static
platform_ptr
create
(
cl_platform_id
platform_id
,
unsigned
start_id
);
private:
platform
(
cl_platform_id
platform_id
,
detail
::
raw_context_ptr
context
,
std
::
string
name
,
std
::
string
vendor
,
std
::
string
version
,
std
::
vector
<
device_ptr
>
devices
);
~
platform
();
static
std
::
string
platform_info
(
cl_platform_id
platform_id
,
unsigned
info_flag
);
cl_platform_id
platform_id_
;
detail
::
raw_context_ptr
context_
;
std
::
string
name_
;
std
::
string
vendor_
;
std
::
string
version_
;
std
::
vector
<
device_ptr
>
devices_
;
};
/******************************************************************************\
* implementation of inline member functions *
\******************************************************************************/
inline
const
std
::
vector
<
device_ptr
>&
platform
::
devices
()
const
{
return
devices_
;
}
inline
const
std
::
string
&
platform
::
name
()
const
{
return
name_
;
}
inline
const
std
::
string
&
platform
::
vendor
()
const
{
return
vendor_
;
}
inline
const
std
::
string
&
platform
::
version
()
const
{
return
version_
;
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/caf/opencl/program.hpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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. *
******************************************************************************/
#pragma once
#include <map>
#include <memory>
#include "caf/ref_counted.hpp"
#include "caf/detail/raw_ptr.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
namespace
caf
{
namespace
opencl
{
class
program
;
using
program_ptr
=
intrusive_ptr
<
program
>
;
/// @brief A wrapper for OpenCL's cl_program.
class
program
:
public
ref_counted
{
public:
friend
class
manager
;
template
<
bool
PassConfig
,
class
...
Ts
>
friend
class
actor_facade
;
template
<
class
T
,
class
...
Ts
>
friend
intrusive_ptr
<
T
>
caf
::
make_counted
(
Ts
&&
...);
private:
program
(
detail
::
raw_context_ptr
context
,
detail
::
raw_command_queue_ptr
queue
,
detail
::
raw_program_ptr
prog
,
std
::
map
<
std
::
string
,
detail
::
raw_kernel_ptr
>
available_kernels
);
~
program
();
detail
::
raw_context_ptr
context_
;
detail
::
raw_program_ptr
program_
;
detail
::
raw_command_queue_ptr
queue_
;
std
::
map
<
std
::
string
,
detail
::
raw_kernel_ptr
>
available_kernels_
;
};
}
// namespace opencl
}
// namespace caf
libcaf_opencl/examples/CMakeLists.txt
deleted
100644 → 0
View file @
b05a4b7b
cmake_minimum_required
(
VERSION 2.8
)
project
(
caf_examples_opencl CXX
)
if
(
OpenCL_LIBRARIES AND NOT CAF_NO_EXAMPLES
)
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
}
${
CAF_EXTRA_LDFLAGS
}
${
CAF_LIBRARIES
}
${
PTHREAD_LIBRARIES
}
${
WSLIB
}
${
OpenCL_LIBRARIES
}
)
add_dependencies
(
${
name
}
all_examples
)
endmacro
()
add
(
proper_matrix .
)
add
(
simple_matrix .
)
add
(
scan .
)
endif
()
libcaf_opencl/examples/proper_matrix.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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/detail/limited_vector.hpp"
#include "caf/opencl/all.hpp"
using
namespace
std
;
using
namespace
caf
;
using
namespace
caf
::
opencl
;
using
caf
::
detail
::
limited_vector
;
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 const float* matrix1,
global const 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
template
<
size_t
Size
>
class
square_matrix
{
public:
using
value_type
=
fvec
::
value_type
;
static
constexpr
size_t
num_elements
=
Size
*
Size
;
// allows serialization
template
<
class
Inspector
>
friend
typename
Inspector
::
result_type
inspect
(
Inspector
&
f
,
square_matrix
&
m
)
{
return
f
(
meta
::
type_name
(
"square_matrix"
),
m
.
data_
);
}
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
()
:
data_
(
num_elements
)
{
}
explicit
square_matrix
(
fvec
d
)
:
data_
(
move
(
d
))
{
assert
(
data_
.
size
()
==
num_elements
);
}
inline
float
&
operator
()(
size_t
column
,
size_t
row
)
{
return
data_
[
column
+
row
*
Size
];
}
inline
const
float
&
operator
()(
size_t
column
,
size_t
row
)
const
{
return
data_
[
column
+
row
*
Size
];
}
inline
void
iota_fill
()
{
iota
(
data_
.
begin
(),
data_
.
end
(),
0
);
}
using
const_iterator
=
typename
fvec
::
const_iterator
;
const_iterator
begin
()
const
{
return
data_
.
begin
();
}
const_iterator
end
()
const
{
return
data_
.
end
();
}
fvec
&
data
()
{
return
data_
;
}
const
fvec
&
data
()
const
{
return
data_
;
}
private:
fvec
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
)
{
auto
&
mngr
=
self
->
system
().
opencl_manager
();
// 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
;
auto
unbox_args
=
[](
message
&
msg
)
->
optional
<
message
>
{
return
msg
.
apply
([](
matrix_type
&
lhs
,
matrix_type
&
rhs
)
{
return
make_message
(
std
::
move
(
lhs
.
data
()),
std
::
move
(
rhs
.
data
()));
});
};
auto
box_res
=
[]
(
fvec
&
result
)
->
message
{
return
make_message
(
matrix_type
{
move
(
result
)});
};
// spawn an opencl actor
// 1st arg: source code of one or more opencl kernels
// 2nd arg: name of the kernel to use
// 3rd arg: the config specifies how many dimensions the kernel uses and how
// many work items are created, creates matrix_size * matrix_size
// global work items in this case
// 4th 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)
// 5th arg: converts the ouptut vector back to a matrix that is then
// used as response message
// from 6 : a description of the kernel signature using in/out/in_out classes
// with the argument type. Since the actor always expects input
// arguments for global memory to be contained in vectors,
// the vector is omitted here.
auto
worker
=
mngr
.
spawn
(
kernel_source
,
kernel_name
,
nd_range
{
dim_vec
{
matrix_size
,
matrix_size
}},
unbox_args
,
box_res
,
in
<
float
>
{},
in
<
float
>
{},
out
<
float
>
{});
// send both matrices to the actor and
// wait for results in form of a matrix_type
self
->
request
(
worker
,
chrono
::
seconds
(
5
),
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
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
()
.
add_message_type
<
fvec
>
(
"float_vector"
)
.
add_message_type
<
matrix_type
>
(
"square_matrix"
);
actor_system
system
{
cfg
};
system
.
spawn
(
multiplier
);
system
.
await_all_actors_done
();
return
0
;
}
libcaf_opencl/examples/scan.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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 <random>
#include <iomanip>
#include <numeric>
#include <iostream>
#include "caf/all.hpp"
#include "caf/opencl/all.hpp"
using
namespace
std
;
using
namespace
caf
;
using
namespace
caf
::
opencl
;
using
caf
::
detail
::
limited_vector
;
namespace
{
using
uval
=
unsigned
;
using
uvec
=
std
::
vector
<
uval
>
;
using
uref
=
mem_ref
<
uval
>
;
constexpr
const
size_t
problem_size
=
23
;
constexpr
const
char
*
kernel_name_1
=
"phase_1"
;
constexpr
const
char
*
kernel_name_2
=
"phase_2"
;
constexpr
const
char
*
kernel_name_3
=
"phase_3"
;
// opencl kernel, exclusive scan
// last parameter is, by convention, the output parameter
constexpr
const
char
*
kernel_source
=
R"__(
/// Global exclusive scan, phase 1. From:
/// - http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html
kernel void phase_1(global uint* restrict data,
global uint* restrict increments,
local uint* tmp, uint len) {
const uint thread = get_local_id(0);
const uint block = get_group_id(0);
const uint threads_per_block = get_local_size(0);
const uint elements_per_block = threads_per_block * 2;
const uint global_offset = block * elements_per_block;
const uint n = elements_per_block;
uint offset = 1;
// A (2 lines) --> load input into shared memory
tmp[2 * thread] = (global_offset + (2 * thread) < len)
? data[global_offset + (2 * thread)] : 0;
tmp[2 * thread + 1] = (global_offset + (2 * thread + 1) < len)
? data[global_offset + (2 * thread + 1)] : 0;
// build sum in place up the tree
for (uint d = n >> 1; d > 0; d >>= 1) {
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// B (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
tmp[bi] += tmp[ai];
}
offset *= 2;
}
// C (2 lines) --> clear the last element
if (thread == 0) {
increments[block] = tmp[n - 1];
tmp[n - 1] = 0;
}
// traverse down tree & build scan
for (uint d = 1; d < n; d *= 2) {
offset >>= 1;
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// D (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
uint t = tmp[ai];
tmp[ai] = tmp[bi];
tmp[bi] += t;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
// E (2 line) --> write results to device memory
if (global_offset + (2 * thread) < len)
data[global_offset + (2 * thread)] = tmp[2 * thread];
if (global_offset + (2 * thread + 1) < len)
data[global_offset + (2 * thread + 1)] = tmp[2 * thread + 1];
}
/// Global exclusive scan, phase 2.
kernel void phase_2(global uint* restrict data, // not used ...
global uint* restrict increments,
uint len) {
local uint tmp[2048];
uint thread = get_local_id(0);
uint offset = 1;
const uint n = 2048;
// A (2 lines) --> load input into shared memory
tmp[2 * thread] = (2 * thread < len) ? increments[2 * thread] : 0;
tmp[2 * thread + 1] = (2 * thread + 1 < len) ? increments[2 * thread + 1] : 0;
// build sum in place up the tree
for (uint d = n >> 1; d > 0; d >>= 1) {
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// B (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
tmp[bi] += tmp[ai];
}
offset *= 2;
}
// C (2 lines) --> clear the last element
if (thread == 0)
tmp[n - 1] = 0;
// traverse down tree & build scan
for (uint d = 1; d < n; d *= 2) {
offset >>= 1;
barrier(CLK_LOCAL_MEM_FENCE);
if (thread < d) {
// D (2 lines)
int ai = offset * (2 * thread + 1) - 1;
int bi = offset * (2 * thread + 2) - 1;
uint t = tmp[ai];
tmp[ai] = tmp[bi];
tmp[bi] += t;
}
}
barrier(CLK_LOCAL_MEM_FENCE);
// E (2 line) --> write results to device memory
if (2 * thread < len) increments[2 * thread] = tmp[2 * thread];
if (2 * thread + 1 < len) increments[2 * thread + 1] = tmp[2 * thread + 1];
}
kernel void phase_3(global uint* restrict data,
global uint* restrict increments,
uint len) {
const uint thread = get_local_id(0);
const uint block = get_group_id(0);
const uint threads_per_block = get_local_size(0);
const uint elements_per_block = threads_per_block * 2;
const uint global_offset = block * elements_per_block;
// add the appropriate value to each block
uint ai = 2 * thread;
uint bi = 2 * thread + 1;
uint ai_global = ai + global_offset;
uint bi_global = bi + global_offset;
uint increment = increments[block];
if (ai_global < len) data[ai_global] += increment;
if (bi_global < len) data[bi_global] += increment;
}
)__"
;
}
// namespace
template
<
class
T
,
class
E
=
caf
::
detail
::
enable_if_t
<
is_integral
<
T
>
::
value
>>
T
round_up
(
T
numToRound
,
T
multiple
)
{
return
((
numToRound
+
multiple
-
1
)
/
multiple
)
*
multiple
;
}
int
main
()
{
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
()
.
add_message_type
<
uvec
>
(
"uint_vector"
);
actor_system
system
{
cfg
};
cout
<<
"Calculating exclusive scan of '"
<<
problem_size
<<
"' values."
<<
endl
;
// ---- create data ----
uvec
values
(
problem_size
);
random_device
rd
;
mt19937
gen
(
rd
());
uniform_int_distribution
<
uval
>
val_gen
(
0
,
1023
);
std
::
generate
(
begin
(
values
),
end
(
values
),
[
&
]()
{
return
val_gen
(
gen
);
});
// ---- find device ----
auto
&
mngr
=
system
.
opencl_manager
();
//
string
prefix
=
"GeForce"
;
auto
opt
=
mngr
.
find_device_if
([
&
](
const
device_ptr
dev
)
{
auto
&
name
=
dev
->
name
();
return
equal
(
begin
(
prefix
),
end
(
prefix
),
begin
(
name
));
});
if
(
!
opt
)
{
cout
<<
"No device starting with '"
<<
prefix
<<
"' found. "
<<
"Will try the first OpenCL device available."
<<
endl
;
opt
=
mngr
.
find_device
();
}
if
(
!
opt
)
{
cerr
<<
"Not OpenCL device available."
<<
endl
;
return
0
;
}
else
{
cerr
<<
"Found device '"
<<
(
*
opt
)
->
name
()
<<
"'."
<<
endl
;
}
{
// ---- general ----
auto
dev
=
move
(
*
opt
);
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
system
};
// ---- config parameters ----
auto
half_block
=
dev
->
max_work_group_size
()
/
2
;
auto
get_size
=
[
half_block
](
size_t
n
)
->
size_t
{
return
round_up
((
n
+
1
)
/
2
,
half_block
);
};
auto
nd_conf
=
[
half_block
,
get_size
](
size_t
dim
)
{
return
nd_range
{
dim_vec
{
get_size
(
dim
)},
{},
dim_vec
{
half_block
}};
};
auto
reduced_ref
=
[
&
](
const
uref
&
,
uval
n
)
{
// calculate number of groups from the group size from the values size
return
size_t
{
get_size
(
n
)
/
half_block
};
};
// default nd-range
auto
ndr
=
nd_range
{
dim_vec
{
half_block
},
{},
dim_vec
{
half_block
}};
// ---- scan actors ----
auto
phase1
=
mngr
.
spawn
(
prog
,
kernel_name_1
,
ndr
,
[
nd_conf
](
nd_range
&
range
,
message
&
msg
)
->
optional
<
message
>
{
return
msg
.
apply
([
&
](
uvec
&
vec
)
{
auto
size
=
vec
.
size
();
range
=
nd_conf
(
size
);
return
make_message
(
std
::
move
(
vec
),
static_cast
<
uval
>
(
size
));
});
},
in_out
<
uval
,
val
,
mref
>
{},
out
<
uval
,
mref
>
{
reduced_ref
},
local
<
uval
>
{
half_block
*
2
},
priv
<
uval
,
val
>
{}
);
auto
phase2
=
mngr
.
spawn
(
prog
,
kernel_name_2
,
ndr
,
[
nd_conf
](
nd_range
&
range
,
message
&
msg
)
->
optional
<
message
>
{
return
msg
.
apply
([
&
](
uref
&
data
,
uref
&
incs
)
{
auto
size
=
incs
.
size
();
range
=
nd_conf
(
size
);
return
make_message
(
move
(
data
),
move
(
incs
),
static_cast
<
uval
>
(
size
));
});
},
in_out
<
uval
,
mref
,
mref
>
{},
in_out
<
uval
,
mref
,
mref
>
{},
priv
<
uval
,
val
>
{}
);
auto
phase3
=
mngr
.
spawn
(
prog
,
kernel_name_3
,
ndr
,
[
nd_conf
](
nd_range
&
range
,
message
&
msg
)
->
optional
<
message
>
{
return
msg
.
apply
([
&
](
uref
&
data
,
uref
&
incs
)
{
auto
size
=
incs
.
size
();
range
=
nd_conf
(
size
);
return
make_message
(
move
(
data
),
move
(
incs
),
static_cast
<
uval
>
(
size
));
});
},
in_out
<
uval
,
mref
,
val
>
{},
in
<
uval
,
mref
>
{},
priv
<
uval
,
val
>
{}
);
// ---- composed scan actor ----
auto
scanner
=
phase3
*
phase2
*
phase1
;
// ---- scan the data ----
self
->
send
(
scanner
,
values
);
self
->
receive
(
[
&
](
const
uvec
&
results
)
{
cout
<<
"Received results."
<<
endl
;
cout
<<
" index | values | scan "
<<
endl
<<
"-------+--------+--------"
<<
endl
;
for
(
size_t
i
=
0
;
i
<
problem_size
;
++
i
)
cout
<<
setw
(
6
)
<<
i
<<
" | "
<<
setw
(
6
)
<<
values
[
i
]
<<
" | "
<<
setw
(
6
)
<<
results
[
i
]
<<
endl
;
}
);
}
system
.
await_all_actors_done
();
return
0
;
}
libcaf_opencl/examples/simple_matrix.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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/all.hpp"
using
namespace
std
;
using
namespace
caf
;
using
namespace
caf
::
opencl
;
using
caf
::
detail
::
limited_vector
;
namespace
{
using
fvec
=
std
::
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 const float* matrix1,
global const 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
void
print_as_matrix
(
const
fvec
&
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
fvec
m1
(
matrix_size
*
matrix_size
);
fvec
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
// 1st arg: source code of one or more kernels
// 2nd arg: name of the kernel to use
// 3rd arg: a spawn configuration that includes:
// - the global dimension arguments for opencl's enqueue
// creates matrix_size * matrix_size global work items
// - offsets for global dimensions (optional)
// - local dimensions (optional)
// 4th to Nth arg: the kernel signature described by in/out/in_out classes
// that contain the argument type in their template. Since the actor
// expects its arguments for global memory to be passed in vectors,
// the vector type is omitted for brevity.
auto
worker
=
self
->
system
().
opencl_manager
().
spawn
(
kernel_source
,
kernel_name
,
nd_range
{
dim_vec
{
matrix_size
,
matrix_size
}},
in
<
float
>
{},
in
<
float
>
{},
out
<
float
>
{}
);
// send both matrices to the actor and wait for a result
self
->
request
(
worker
,
chrono
::
seconds
(
5
),
move
(
m1
),
move
(
m2
)).
then
(
[](
const
fvec
&
result
)
{
cout
<<
"result: "
<<
endl
;
print_as_matrix
(
result
);
}
);
}
int
main
()
{
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
()
.
add_message_type
<
fvec
>
(
"float_vector"
);
actor_system
system
{
cfg
};
system
.
spawn
(
multiplier
);
system
.
await_all_actors_done
();
return
0
;
}
libcaf_opencl/src/device.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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 <iostream>
#include <utility>
#include "caf/logger.hpp"
#include "caf/ref_counted.hpp"
#include "caf/string_algorithms.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/opencl_err.hpp"
using
namespace
std
;
namespace
caf
{
namespace
opencl
{
device_ptr
device
::
create
(
const
detail
::
raw_context_ptr
&
context
,
const
detail
::
raw_device_ptr
&
device_id
,
unsigned
id
)
{
CAF_LOG_DEBUG
(
"creating device for opencl device with id:"
<<
CAF_ARG
(
id
));
// look up properties we need to create the command queue
auto
supported
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_QUEUE_PROPERTIES
);
bool
profiling
=
false
;
// (supported & CL_QUEUE_PROFILING_ENABLE) != 0u;
bool
out_of_order
=
(
supported
&
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
)
!=
0u
;
unsigned
properties
=
profiling
?
CL_QUEUE_PROFILING_ENABLE
:
0
;
if
(
out_of_order
)
properties
|=
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
;
// create the command queue
detail
::
raw_command_queue_ptr
command_queue
{
v2get
(
CAF_CLF
(
clCreateCommandQueue
),
context
.
get
(),
device_id
.
get
(),
properties
),
false
};
// create the device
auto
dev
=
make_counted
<
device
>
(
device_id
,
std
::
move
(
command_queue
),
context
,
id
);
//device dev{device_id, std::move(command_queue), context, id};
// look up device properties
dev
->
address_bits_
=
info
<
cl_uint
>
(
device_id
,
CL_DEVICE_ADDRESS_BITS
);
dev
->
little_endian_
=
info
<
cl_bool
>
(
device_id
,
CL_DEVICE_ENDIAN_LITTLE
);
dev
->
global_mem_cache_size_
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
);
dev
->
global_mem_cacheline_size_
=
info
<
cl_uint
>
(
device_id
,
CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
);
dev
->
global_mem_size_
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_GLOBAL_MEM_SIZE
);
dev
->
host_unified_memory_
=
info
<
cl_bool
>
(
device_id
,
CL_DEVICE_HOST_UNIFIED_MEMORY
);
dev
->
local_mem_size_
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_LOCAL_MEM_SIZE
);
dev
->
local_mem_type_
=
info
<
cl_uint
>
(
device_id
,
CL_DEVICE_LOCAL_MEM_TYPE
);
dev
->
max_clock_frequency_
=
info
<
cl_uint
>
(
device_id
,
CL_DEVICE_MAX_CLOCK_FREQUENCY
);
dev
->
max_compute_units_
=
info
<
cl_uint
>
(
device_id
,
CL_DEVICE_MAX_COMPUTE_UNITS
);
dev
->
max_constant_args_
=
info
<
cl_uint
>
(
device_id
,
CL_DEVICE_MAX_CONSTANT_ARGS
);
dev
->
max_constant_buffer_size_
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
);
dev
->
max_mem_alloc_size_
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_MAX_MEM_ALLOC_SIZE
);
dev
->
max_parameter_size_
=
info
<
size_t
>
(
device_id
,
CL_DEVICE_MAX_PARAMETER_SIZE
);
dev
->
max_work_group_size_
=
info
<
size_t
>
(
device_id
,
CL_DEVICE_MAX_WORK_GROUP_SIZE
);
dev
->
max_work_item_dimensions_
=
info
<
cl_uint
>
(
device_id
,
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
);
dev
->
profiling_timer_resolution_
=
info
<
size_t
>
(
device_id
,
CL_DEVICE_PROFILING_TIMER_RESOLUTION
);
dev
->
max_work_item_sizes_
.
resize
(
dev
->
max_work_item_dimensions_
);
clGetDeviceInfo
(
device_id
.
get
(),
CL_DEVICE_MAX_WORK_ITEM_SIZES
,
sizeof
(
size_t
)
*
dev
->
max_work_item_dimensions_
,
dev
->
max_work_item_sizes_
.
data
(),
nullptr
);
dev
->
device_type_
=
device_type_from_ulong
(
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_TYPE
));
string
extensions
=
info_string
(
device_id
,
CL_DEVICE_EXTENSIONS
);
split
(
dev
->
extensions_
,
extensions
,
" "
,
false
);
dev
->
opencl_c_version_
=
info_string
(
device_id
,
CL_DEVICE_EXTENSIONS
);
dev
->
device_vendor_
=
info_string
(
device_id
,
CL_DEVICE_VENDOR
);
dev
->
device_version_
=
info_string
(
device_id
,
CL_DEVICE_VERSION
);
dev
->
name_
=
info_string
(
device_id
,
CL_DEVICE_NAME
);
return
dev
;
}
void
device
::
synchronize
()
{
clFinish
(
queue_
.
get
());
}
string
device
::
info_string
(
const
detail
::
raw_device_ptr
&
device_id
,
unsigned
info_flag
)
{
size_t
size
;
clGetDeviceInfo
(
device_id
.
get
(),
info_flag
,
0
,
nullptr
,
&
size
);
vector
<
char
>
buffer
(
size
);
clGetDeviceInfo
(
device_id
.
get
(),
info_flag
,
sizeof
(
char
)
*
size
,
buffer
.
data
(),
nullptr
);
return
string
(
buffer
.
data
());
}
device
::
device
(
detail
::
raw_device_ptr
device_id
,
detail
::
raw_command_queue_ptr
queue
,
detail
::
raw_context_ptr
context
,
unsigned
id
)
:
device_id_
(
std
::
move
(
device_id
)),
queue_
(
std
::
move
(
queue
)),
context_
(
std
::
move
(
context
)),
id_
(
id
)
{
// nop
}
device
::~
device
()
{
// nop
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/src/global.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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 <string>
#include <sstream>
#include "caf/opencl/global.hpp"
cl_int
clReleaseDeviceDummy
(
cl_device_id
)
{
return
0
;
}
cl_int
clRetainDeviceDummy
(
cl_device_id
)
{
return
0
;
}
namespace
caf
{
namespace
opencl
{
std
::
ostream
&
operator
<<
(
std
::
ostream
&
os
,
device_type
dev
)
{
switch
(
dev
)
{
case
def
:
os
<<
"default"
;
break
;
case
cpu
:
os
<<
"CPU"
;
break
;
case
gpu
:
os
<<
"GPU"
;
break
;
case
accelerator
:
os
<<
"accelerator"
;
break
;
case
custom
:
os
<<
"custom"
;
break
;
case
all
:
os
<<
"all"
;
break
;
default
:
os
.
setstate
(
std
::
ios_base
::
failbit
);
}
return
os
;
}
device_type
device_type_from_ulong
(
cl_ulong
dev
)
{
switch
(
dev
)
{
case
CL_DEVICE_TYPE_CPU
:
return
cpu
;
case
CL_DEVICE_TYPE_GPU
:
return
gpu
;
case
CL_DEVICE_TYPE_ACCELERATOR
:
return
accelerator
;
case
CL_DEVICE_TYPE_CUSTOM
:
return
custom
;
case
CL_DEVICE_TYPE_ALL
:
return
all
;
default
:
return
def
;
}
}
std
::
string
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"
;
// error codes used by extensions
// see: http://streamcomputing.eu/blog/2013-04-28/opencl-1-2-error-codes/
case
-
1000
:
return
"CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR"
;
case
-
1001
:
return
"No valid ICDs found"
;
case
-
1002
:
return
"CL_INVALID_D3D10_DEVICE_KHR"
;
case
-
1003
:
return
"CL_INVALID_D3D10_RESOURCE_KHR"
;
case
-
1004
:
return
"CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR"
;
case
-
1005
:
return
"CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR"
;
default:
return
"UNKNOWN_ERROR: "
+
std
::
to_string
(
err
);
}
}
std
::
string
event_status
(
cl_event
e
)
{
std
::
stringstream
ss
;
cl_int
s
;
auto
err
=
clGetEventInfo
(
e
,
CL_EVENT_COMMAND_EXECUTION_STATUS
,
sizeof
(
cl_int
),
&
s
,
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
ss
<<
std
::
string
(
"ERROR "
)
+
std
::
to_string
(
s
);
return
ss
.
str
();
}
cl_command_type
t
;
err
=
clGetEventInfo
(
e
,
CL_EVENT_COMMAND_TYPE
,
sizeof
(
cl_command_type
),
&
t
,
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
ss
<<
std
::
string
(
"ERROR "
)
+
std
::
to_string
(
s
);
return
ss
.
str
();
}
switch
(
s
)
{
case
(
CL_QUEUED
):
ss
<<
std
::
string
(
"CL_QUEUED"
);
break
;
case
(
CL_SUBMITTED
):
ss
<<
std
::
string
(
"CL_SUBMITTED"
);
break
;
case
(
CL_RUNNING
):
ss
<<
std
::
string
(
"CL_RUNNING"
);
break
;
case
(
CL_COMPLETE
):
ss
<<
std
::
string
(
"CL_COMPLETE"
);
break
;
default:
ss
<<
std
::
string
(
"DEFAULT "
)
+
std
::
to_string
(
s
);
return
ss
.
str
();
}
ss
<<
" / "
;
switch
(
t
)
{
case
(
CL_COMMAND_NDRANGE_KERNEL
):
ss
<<
std
::
string
(
"CL_COMMAND_NDRANGE_KERNEL"
);
break
;
case
(
CL_COMMAND_TASK
):
ss
<<
std
::
string
(
"CL_COMMAND_TASK"
);
break
;
case
(
CL_COMMAND_NATIVE_KERNEL
):
ss
<<
std
::
string
(
"CL_COMMAND_NATIVE_KERNEL"
);
break
;
case
(
CL_COMMAND_READ_BUFFER
):
ss
<<
std
::
string
(
"CL_COMMAND_READ_BUFFER"
);
break
;
case
(
CL_COMMAND_WRITE_BUFFER
):
ss
<<
std
::
string
(
"CL_COMMAND_WRITE_BUFFER"
);
break
;
case
(
CL_COMMAND_COPY_BUFFER
):
ss
<<
std
::
string
(
"CL_COMMAND_COPY_BUFFER"
);
break
;
case
(
CL_COMMAND_READ_IMAGE
):
ss
<<
std
::
string
(
"CL_COMMAND_READ_IMAGE"
);
break
;
case
(
CL_COMMAND_WRITE_IMAGE
):
ss
<<
std
::
string
(
"CL_COMMAND_WRITE_IMAGE"
);
break
;
case
(
CL_COMMAND_COPY_IMAGE
):
ss
<<
std
::
string
(
"CL_COMMAND_COPY_IMAGE"
);
break
;
case
(
CL_COMMAND_COPY_BUFFER_TO_IMAGE
):
ss
<<
std
::
string
(
"CL_COMMAND_COPY_BUFFER_TO_IMAGE"
);
break
;
case
(
CL_COMMAND_COPY_IMAGE_TO_BUFFER
):
ss
<<
std
::
string
(
"CL_COMMAND_COPY_IMAGE_TO_BUFFER"
);
break
;
case
(
CL_COMMAND_MAP_BUFFER
):
ss
<<
std
::
string
(
"CL_COMMAND_MAP_BUFFER"
);
break
;
case
(
CL_COMMAND_MAP_IMAGE
):
ss
<<
std
::
string
(
"CL_COMMAND_MAP_IMAGE"
);
break
;
case
(
CL_COMMAND_UNMAP_MEM_OBJECT
):
ss
<<
std
::
string
(
"CL_COMMAND_UNMAP_MEM_OBJECT"
);
break
;
case
(
CL_COMMAND_MARKER
):
ss
<<
std
::
string
(
"CL_COMMAND_MARKER"
);
break
;
case
(
CL_COMMAND_ACQUIRE_GL_OBJECTS
):
ss
<<
std
::
string
(
"CL_COMMAND_ACQUIRE_GL_OBJECTS"
);
break
;
case
(
CL_COMMAND_RELEASE_GL_OBJECTS
):
ss
<<
std
::
string
(
"CL_COMMAND_RELEASE_GL_OBJECTS"
);
break
;
case
(
CL_COMMAND_READ_BUFFER_RECT
):
ss
<<
std
::
string
(
"CL_COMMAND_READ_BUFFER_RECT"
);
break
;
case
(
CL_COMMAND_WRITE_BUFFER_RECT
):
ss
<<
std
::
string
(
"CL_COMMAND_WRITE_BUFFER_RECT"
);
break
;
case
(
CL_COMMAND_COPY_BUFFER_RECT
):
ss
<<
std
::
string
(
"CL_COMMAND_COPY_BUFFER_RECT"
);
break
;
case
(
CL_COMMAND_USER
):
ss
<<
std
::
string
(
"CL_COMMAND_USER"
);
break
;
case
(
CL_COMMAND_BARRIER
):
ss
<<
std
::
string
(
"CL_COMMAND_BARRIER"
);
break
;
case
(
CL_COMMAND_MIGRATE_MEM_OBJECTS
):
ss
<<
std
::
string
(
"CL_COMMAND_MIGRATE_MEM_OBJECTS"
);
break
;
case
(
CL_COMMAND_FILL_BUFFER
):
ss
<<
std
::
string
(
"CL_COMMAND_FILL_BUFFER"
);
break
;
case
(
CL_COMMAND_FILL_IMAGE
):
ss
<<
std
::
string
(
"CL_COMMAND_FILL_IMAGE"
);
break
;
default:
ss
<<
std
::
string
(
"DEFAULT "
)
+
std
::
to_string
(
s
);
return
ss
.
str
();
}
return
ss
.
str
();
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/src/manager.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Raphael Hiesgen <raphael.hiesgen (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 LICENSE_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 <fstream>
#include "caf/detail/type_list.hpp"
#include "caf/raise_error.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/manager.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
using
namespace
std
;
namespace
caf
{
namespace
opencl
{
optional
<
device_ptr
>
manager
::
find_device
(
size_t
dev_id
)
const
{
if
(
platforms_
.
empty
())
return
none
;
size_t
to
=
0
;
for
(
auto
&
pl
:
platforms_
)
{
auto
from
=
to
;
to
+=
pl
->
devices
().
size
();
if
(
dev_id
>=
from
&&
dev_id
<
to
)
return
pl
->
devices
()[
dev_id
-
from
];
}
return
none
;
}
void
manager
::
init
(
actor_system_config
&
)
{
// get number of available platforms
auto
num_platforms
=
v1get
<
cl_uint
>
(
CAF_CLF
(
clGetPlatformIDs
));
// get platform ids
std
::
vector
<
cl_platform_id
>
platform_ids
(
num_platforms
);
v2callcl
(
CAF_CLF
(
clGetPlatformIDs
),
num_platforms
,
platform_ids
.
data
());
if
(
platform_ids
.
empty
())
CAF_RAISE_ERROR
(
"no OpenCL platform found"
);
// initialize platforms (device discovery)
unsigned
current_device_id
=
0
;
for
(
auto
&
pl_id
:
platform_ids
)
{
platforms_
.
push_back
(
platform
::
create
(
pl_id
,
current_device_id
));
current_device_id
+=
static_cast
<
unsigned
>
(
platforms_
.
back
()
->
devices
().
size
());
}
}
void
manager
::
start
()
{
// nop
}
void
manager
::
stop
()
{
// nop
}
actor_system
::
module
::
id_t
manager
::
id
()
const
{
return
actor_system
::
module
::
opencl_manager
;
}
void
*
manager
::
subtype_ptr
()
{
return
this
;
}
actor_system
::
module
*
manager
::
make
(
actor_system
&
sys
,
caf
::
detail
::
type_list
<>
)
{
return
new
manager
{
sys
};
}
program_ptr
manager
::
create_program_from_file
(
const
char
*
path
,
const
char
*
options
,
uint32_t
device_id
)
{
std
::
ifstream
read_source
{
std
::
string
(
path
),
std
::
ios
::
in
};
string
kernel_source
;
if
(
read_source
)
{
read_source
.
seekg
(
0
,
std
::
ios
::
end
);
kernel_source
.
resize
(
static_cast
<
size_t
>
(
read_source
.
tellg
()));
read_source
.
seekg
(
0
,
std
::
ios
::
beg
);
read_source
.
read
(
&
kernel_source
[
0
],
static_cast
<
streamsize
>
(
kernel_source
.
size
()));
read_source
.
close
();
}
else
{
CAF_RAISE_ERROR
(
"create_program_from_file: path not found"
);
}
return
create_program
(
kernel_source
.
c_str
(),
options
,
device_id
);
}
program_ptr
manager
::
create_program
(
const
char
*
kernel_source
,
const
char
*
options
,
uint32_t
device_id
)
{
auto
dev
=
find_device
(
device_id
);
if
(
!
dev
)
{
CAF_RAISE_ERROR
(
"create_program: no device found"
);
}
return
create_program
(
kernel_source
,
options
,
*
dev
);
}
program_ptr
manager
::
create_program_from_file
(
const
char
*
path
,
const
char
*
options
,
const
device_ptr
dev
)
{
std
::
ifstream
read_source
{
std
::
string
(
path
),
std
::
ios
::
in
};
string
kernel_source
;
if
(
read_source
)
{
read_source
.
seekg
(
0
,
std
::
ios
::
end
);
kernel_source
.
resize
(
static_cast
<
size_t
>
(
read_source
.
tellg
()));
read_source
.
seekg
(
0
,
std
::
ios
::
beg
);
read_source
.
read
(
&
kernel_source
[
0
],
static_cast
<
streamsize
>
(
kernel_source
.
size
()));
read_source
.
close
();
}
else
{
CAF_RAISE_ERROR
(
"create_program_from_file: path not found"
);
}
return
create_program
(
kernel_source
.
c_str
(),
options
,
dev
);
}
program_ptr
manager
::
create_program
(
const
char
*
kernel_source
,
const
char
*
options
,
const
device_ptr
dev
)
{
// create program object from kernel source
size_t
kernel_source_length
=
strlen
(
kernel_source
);
detail
::
raw_program_ptr
pptr
;
pptr
.
reset
(
v2get
(
CAF_CLF
(
clCreateProgramWithSource
),
dev
->
context_
.
get
(),
1u
,
&
kernel_source
,
&
kernel_source_length
),
false
);
// build programm from program object
auto
dev_tmp
=
dev
->
device_id_
.
get
();
auto
err
=
clBuildProgram
(
pptr
.
get
(),
1
,
&
dev_tmp
,
options
,
nullptr
,
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
==
CL_BUILD_PROGRAM_FAILURE
)
{
size_t
buildlog_buffer_size
=
0
;
// get the log length
clGetProgramBuildInfo
(
pptr
.
get
(),
dev_tmp
,
CL_PROGRAM_BUILD_LOG
,
0
,
nullptr
,
&
buildlog_buffer_size
);
vector
<
char
>
buffer
(
buildlog_buffer_size
);
// fill the buffer with buildlog informations
clGetProgramBuildInfo
(
pptr
.
get
(),
dev_tmp
,
CL_PROGRAM_BUILD_LOG
,
sizeof
(
char
)
*
buildlog_buffer_size
,
buffer
.
data
(),
nullptr
);
ostringstream
ss
;
ss
<<
"############## Build log ##############"
<<
endl
<<
string
(
buffer
.
data
())
<<
endl
<<
"#######################################"
;
// seems that just apple implemented the
// pfn_notify callback, but we can get
// the build log
#ifndef CAF_MACOS
CAF_LOG_ERROR
(
CAF_ARG
(
ss
.
str
()));
#endif
}
CAF_RAISE_ERROR
(
"clBuildProgram failed"
);
}
cl_uint
number_of_kernels
=
0
;
clCreateKernelsInProgram
(
pptr
.
get
(),
0u
,
nullptr
,
&
number_of_kernels
);
map
<
string
,
detail
::
raw_kernel_ptr
>
available_kernels
;
if
(
number_of_kernels
>
0
)
{
vector
<
cl_kernel
>
kernels
(
number_of_kernels
);
err
=
clCreateKernelsInProgram
(
pptr
.
get
(),
number_of_kernels
,
kernels
.
data
(),
nullptr
);
if
(
err
!=
CL_SUCCESS
)
CAF_RAISE_ERROR
(
"clCreateKernelsInProgram failed"
);
for
(
cl_uint
i
=
0
;
i
<
number_of_kernels
;
++
i
)
{
size_t
len
;
clGetKernelInfo
(
kernels
[
i
],
CL_KERNEL_FUNCTION_NAME
,
0
,
nullptr
,
&
len
);
vector
<
char
>
name
(
len
);
err
=
clGetKernelInfo
(
kernels
[
i
],
CL_KERNEL_FUNCTION_NAME
,
len
,
reinterpret_cast
<
void
*>
(
name
.
data
()),
nullptr
);
if
(
err
!=
CL_SUCCESS
)
CAF_RAISE_ERROR
(
"clGetKernelInfo failed"
);
detail
::
raw_kernel_ptr
kernel
;
kernel
.
reset
(
move
(
kernels
[
i
]));
available_kernels
.
emplace
(
string
(
name
.
data
()),
move
(
kernel
));
}
}
else
{
CAF_LOG_WARNING
(
"Could not built all kernels in program. Since this happens"
" on some platforms, we'll ignore this and try to build"
" each kernel individually by name."
);
}
return
make_counted
<
program
>
(
dev
->
context_
,
dev
->
queue_
,
pptr
,
move
(
available_kernels
));
}
manager
::
manager
(
actor_system
&
sys
)
:
system_
(
sys
)
{
// nop
}
manager
::~
manager
()
{
// nop
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/src/opencl_err.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* Raphael Hiesgen <raphael.hiesgen (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 LICENSE_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 "caf/opencl/opencl_err.hpp"
#include "caf/logger.hpp"
namespace
caf
{
namespace
opencl
{
void
throwcl
(
const
char
*
,
cl_int
err
)
{
if
(
err
!=
CL_SUCCESS
)
{
CAF_RAISE_ERROR
(
"throwcl: unrecoverable OpenCL error"
);
}
}
void
CL_CALLBACK
pfn_notify
(
const
char
*
errinfo
,
const
void
*
,
size_t
,
void
*
)
{
CAF_LOG_ERROR
(
"
\n
##### Error message via pfn_notify #####
\n
"
<<
errinfo
<<
"
\n
########################################"
);
CAF_IGNORE_UNUSED
(
errinfo
);
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/src/platform.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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 <utility>
#include <vector>
#include <iostream>
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
using
namespace
std
;
namespace
caf
{
namespace
opencl
{
platform_ptr
platform
::
create
(
cl_platform_id
platform_id
,
unsigned
start_id
)
{
vector
<
unsigned
>
device_types
=
{
CL_DEVICE_TYPE_GPU
,
CL_DEVICE_TYPE_ACCELERATOR
,
CL_DEVICE_TYPE_CPU
};
vector
<
cl_device_id
>
ids
;
for
(
cl_device_type
device_type
:
device_types
)
{
auto
known
=
ids
.
size
();
cl_uint
discoverd
;
auto
err
=
clGetDeviceIDs
(
platform_id
,
device_type
,
0
,
nullptr
,
&
discoverd
);
if
(
err
==
CL_DEVICE_NOT_FOUND
)
{
continue
;
// no devices of the type found
}
else
if
(
err
!=
CL_SUCCESS
)
{
throwcl
(
"clGetDeviceIDs"
,
err
);
}
ids
.
resize
(
known
+
discoverd
);
v2callcl
(
CAF_CLF
(
clGetDeviceIDs
),
platform_id
,
device_type
,
discoverd
,
(
ids
.
data
()
+
known
));
}
vector
<
detail
::
raw_device_ptr
>
devices
;
devices
.
resize
(
ids
.
size
());
auto
lift
=
[](
cl_device_id
ptr
)
{
return
detail
::
raw_device_ptr
{
ptr
,
false
};
};
transform
(
ids
.
begin
(),
ids
.
end
(),
devices
.
begin
(),
lift
);
detail
::
raw_context_ptr
context
;
context
.
reset
(
v2get
(
CAF_CLF
(
clCreateContext
),
nullptr
,
static_cast
<
unsigned
>
(
ids
.
size
()),
ids
.
data
(),
pfn_notify
,
nullptr
),
false
);
vector
<
device_ptr
>
device_information
;
for
(
auto
&
device_id
:
devices
)
{
device_information
.
push_back
(
device
::
create
(
context
,
device_id
,
start_id
++
));
}
if
(
device_information
.
empty
())
CAF_RAISE_ERROR
(
"no devices for the platform found"
);
auto
name
=
platform_info
(
platform_id
,
CL_PLATFORM_NAME
);
auto
vendor
=
platform_info
(
platform_id
,
CL_PLATFORM_VENDOR
);
auto
version
=
platform_info
(
platform_id
,
CL_PLATFORM_VERSION
);
return
make_counted
<
platform
>
(
platform_id
,
move
(
context
),
move
(
name
),
move
(
vendor
),
move
(
version
),
move
(
device_information
));
}
string
platform
::
platform_info
(
cl_platform_id
platform_id
,
unsigned
info_flag
)
{
size_t
size
;
auto
err
=
clGetPlatformInfo
(
platform_id
,
info_flag
,
0
,
nullptr
,
&
size
);
throwcl
(
"clGetPlatformInfo"
,
err
);
vector
<
char
>
buffer
(
size
);
v2callcl
(
CAF_CLF
(
clGetPlatformInfo
),
platform_id
,
info_flag
,
sizeof
(
char
)
*
size
,
buffer
.
data
());
return
string
(
buffer
.
data
());
}
platform
::
platform
(
cl_platform_id
platform_id
,
detail
::
raw_context_ptr
context
,
string
name
,
string
vendor
,
string
version
,
vector
<
device_ptr
>
devices
)
:
platform_id_
(
platform_id
),
context_
(
std
::
move
(
context
)),
name_
(
move
(
name
)),
vendor_
(
move
(
vendor
)),
version_
(
move
(
version
)),
devices_
(
move
(
devices
))
{
// nop
}
platform
::~
platform
()
{
// nop
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/src/program.cpp
deleted
100644 → 0
View file @
b05a4b7b
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* *
* 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 LICENSE_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 <map>
#include <vector>
#include <string>
#include <cstring>
#include <iostream>
#include "caf/opencl/manager.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/opencl_err.hpp"
using
namespace
std
;
namespace
caf
{
namespace
opencl
{
program
::
program
(
detail
::
raw_context_ptr
context
,
detail
::
raw_command_queue_ptr
queue
,
detail
::
raw_program_ptr
prog
,
map
<
string
,
detail
::
raw_kernel_ptr
>
available_kernels
)
:
context_
(
move
(
context
)),
program_
(
move
(
prog
)),
queue_
(
move
(
queue
)),
available_kernels_
(
move
(
available_kernels
))
{
// nop
}
program
::~
program
()
{
// nop
}
}
// namespace opencl
}
// namespace caf
libcaf_opencl/test/opencl.cpp
deleted
100644 → 0
View file @
b05a4b7b
#define CAF_SUITE opencl
#include "caf/test/unit_test.hpp"
#include <vector>
#include <iomanip>
#include <cassert>
#include <iostream>
#include <algorithm>
#include "caf/all.hpp"
#include "caf/system_messages.hpp"
#include "caf/opencl/all.hpp"
using
namespace
std
;
using
namespace
caf
;
using
namespace
caf
::
opencl
;
using
caf
::
detail
::
tl_at
;
using
caf
::
detail
::
tl_head
;
using
caf
::
detail
::
type_list
;
using
caf
::
detail
::
limited_vector
;
namespace
{
using
ivec
=
vector
<
int
>
;
using
iref
=
mem_ref
<
int
>
;
using
dims
=
opencl
::
dim_vec
;
constexpr
size_t
matrix_size
=
4
;
constexpr
size_t
array_size
=
32
;
constexpr
size_t
problem_size
=
1024
;
constexpr
const
char
*
kn_matrix
=
"matrix_square"
;
constexpr
const
char
*
kn_compiler_flag
=
"compiler_flag"
;
constexpr
const
char
*
kn_reduce
=
"reduce"
;
constexpr
const
char
*
kn_const
=
"const_mod"
;
constexpr
const
char
*
kn_inout
=
"times_two"
;
constexpr
const
char
*
kn_scratch
=
"use_scratch"
;
constexpr
const
char
*
kn_local
=
"use_local"
;
constexpr
const
char
*
kn_order
=
"test_order"
;
constexpr
const
char
*
kn_private
=
"use_private"
;
constexpr
const
char
*
kn_varying
=
"varying"
;
constexpr
const
char
*
compiler_flag
=
"-D CAF_OPENCL_TEST_FLAG"
;
constexpr
const
char
*
kernel_source
=
R"__(
kernel void matrix_square(global const int* restrict matrix,
global int* restrict output) {
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);
int result = 0;
for (size_t idx = 0; idx < size; ++idx) {
result += matrix[idx + y * size] * matrix[x + idx * size];
}
output[x + y * size] = result;
}
// http://developer.amd.com/resources/documentation-articles/
// articles-whitepapers/opencl-optimization-case-study-simple-reductions
kernel void reduce(global const int* restrict buffer,
global int* restrict result) {
local int scratch[512];
int local_index = get_local_id(0);
scratch[local_index] = buffer[get_global_id(0)];
barrier(CLK_LOCAL_MEM_FENCE);
for(int offset = get_local_size(0) / 2; offset > 0; offset = offset / 2) {
if (local_index < offset) {
int other = scratch[local_index + offset];
int mine = scratch[local_index];
scratch[local_index] = (mine < other) ? mine : other;
}
barrier(CLK_LOCAL_MEM_FENCE);
}
if (local_index == 0)
result[get_group_id(0)] = scratch[0];
}
kernel void const_mod(constant int* restrict input,
global int* restrict output) {
size_t idx = get_global_id(0);
output[idx] = input[0];
}
kernel void times_two(global int* restrict values) {
size_t idx = get_global_id(0);
values[idx] = values[idx] * 2;
}
kernel void use_scratch(global int* restrict values,
global int* restrict buf) {
size_t idx = get_global_id(0);
buf[idx] = values[idx];
buf[idx] += values[idx];
values[idx] = buf[idx];
}
inline void prefix_sum(local int* restrict data, size_t len, size_t lids) {
size_t lid = get_local_id(0);
size_t inc = 2;
// reduce
while (inc <= len) {
int j = inc >> 1;
for (int i = (j - 1) + (lid * inc); (i + inc) < len; i += (lids * inc))
data[i + j] = data[i] + data[i + j];
inc = inc << 1;
barrier(CLK_LOCAL_MEM_FENCE);
}
// downsweep
data[len - 1] = 0;
barrier(CLK_LOCAL_MEM_FENCE);
while (inc >= 2) {
int j = inc >> 1;
for (int i = (j - 1) + (lid * inc); (i + j) <= len; i += (lids * inc)) {
uint tmp = data[i + j];
data[i + j] = data[i] + data[i + j];
data[i] = tmp;
}
inc = inc >> 1;
barrier(CLK_LOCAL_MEM_FENCE);
}
}
kernel void use_local(global int* restrict values,
local int* restrict buf) {
size_t lid = get_local_id(0);
size_t gid = get_group_id(0);
size_t gs = get_local_size(0);
buf[lid] = values[gid * gs + lid];
barrier(CLK_LOCAL_MEM_FENCE);
prefix_sum(buf, gs, gs);
barrier(CLK_LOCAL_MEM_FENCE);
values[gid * gs + lid] = buf[lid];
}
kernel void test_order(local int* buf,
global int* restrict values) {
size_t lid = get_local_id(0);
size_t gid = get_group_id(0);
size_t gs = get_local_size(0);
buf[lid] = values[gid * gs + lid];
barrier(CLK_LOCAL_MEM_FENCE);
prefix_sum(buf, gs, gs);
barrier(CLK_LOCAL_MEM_FENCE);
values[gid * gs + lid] = buf[lid];
}
kernel void use_private(global int* restrict buf,
private int val) {
buf[get_global_id(0)] += val;
}
kernel void varying(global const int* restrict in1,
global int* restrict out1,
global const int* restrict in2,
global int* restrict out2) {
size_t idx = get_global_id(0);
out1[idx] = in1[idx];
out2[idx] = in2[idx];
}
)__"
;
#ifndef CAF_NO_EXCEPTIONS
constexpr
const
char
*
kernel_source_error
=
R"__(
kernel void missing(global int*) {
size_t semicolon_missing
}
)__"
;
#endif // CAF_NO_EXCEPTIONS
constexpr
const
char
*
kernel_source_compiler_flag
=
R"__(
kernel void compiler_flag(global const int* restrict input,
global int* restrict output) {
size_t x = get_global_id(0);
# ifdef CAF_OPENCL_TEST_FLAG
output[x] = input[x];
# else
output[x] = 0;
# endif
}
)__"
;
}
// namespace
template
<
size_t
Size
>
class
square_matrix
{
public:
using
value_type
=
ivec
::
value_type
;
static
constexpr
size_t
num_elements
=
Size
*
Size
;
template
<
class
Inspector
>
friend
typename
Inspector
::
result_type
inspect
(
Inspector
&
f
,
square_matrix
&
x
)
{
return
f
(
meta
::
type_name
(
"square_matrix"
),
x
.
data_
);
}
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
()
:
data_
(
num_elements
)
{
// nop
}
explicit
square_matrix
(
ivec
d
)
:
data_
(
move
(
d
))
{
assert
(
data_
.
size
()
==
num_elements
);
}
int
&
operator
()(
size_t
column
,
size_t
row
)
{
return
data_
[
column
+
row
*
Size
];
}
const
int
&
operator
()(
size_t
column
,
size_t
row
)
const
{
return
data_
[
column
+
row
*
Size
];
}
using
const_iterator
=
typename
ivec
::
const_iterator
;
const_iterator
begin
()
const
{
return
data_
.
begin
();
}
const_iterator
end
()
const
{
return
data_
.
end
();
}
ivec
&
data
()
{
return
data_
;
}
const
ivec
&
data
()
const
{
return
data_
;
}
void
data
(
ivec
new_data
)
{
data_
=
move
(
new_data
);
}
private:
ivec
data_
;
};
template
<
class
T
>
vector
<
T
>
make_iota_vector
(
size_t
num_elements
)
{
vector
<
T
>
result
;
result
.
resize
(
num_elements
);
iota
(
result
.
begin
(),
result
.
end
(),
T
{
0
});
return
result
;
}
template
<
size_t
Size
>
square_matrix
<
Size
>
make_iota_matrix
()
{
square_matrix
<
Size
>
result
;
iota
(
result
.
data
().
begin
(),
result
.
data
().
end
(),
0
);
return
result
;
}
template
<
size_t
Size
>
bool
operator
==
(
const
square_matrix
<
Size
>&
lhs
,
const
square_matrix
<
Size
>&
rhs
)
{
return
lhs
.
data
()
==
rhs
.
data
();
}
template
<
size_t
Size
>
bool
operator
!=
(
const
square_matrix
<
Size
>&
lhs
,
const
square_matrix
<
Size
>&
rhs
)
{
return
!
(
lhs
==
rhs
);
}
using
matrix_type
=
square_matrix
<
matrix_size
>
;
template
<
class
T
>
void
check_vector_results
(
const
string
&
description
,
const
vector
<
T
>&
expected
,
const
vector
<
T
>&
result
)
{
auto
cond
=
(
expected
==
result
);
CAF_CHECK
(
cond
);
if
(
!
cond
)
{
CAF_ERROR
(
description
<<
" failed."
);
cout
<<
"Expected: "
<<
endl
;
for
(
size_t
i
=
0
;
i
<
expected
.
size
();
++
i
)
{
cout
<<
" "
<<
expected
[
i
];
}
cout
<<
endl
<<
"Received: "
<<
endl
;
for
(
size_t
i
=
0
;
i
<
result
.
size
();
++
i
)
{
cout
<<
" "
<<
result
[
i
];
}
cout
<<
endl
;
cout
<<
"Size: "
<<
expected
.
size
()
<<
" vs. "
<<
result
.
size
()
<<
endl
;
cout
<<
"Differ at: "
<<
endl
;
bool
same
=
true
;
for
(
size_t
i
=
0
;
i
<
min
(
expected
.
size
(),
result
.
size
());
++
i
)
{
if
(
expected
[
i
]
!=
result
[
i
])
{
cout
<<
"["
<<
i
<<
"] "
<<
expected
[
i
]
<<
" != "
<<
result
[
i
]
<<
endl
;
same
=
false
;
}
}
if
(
same
)
{
cout
<<
"... nowhere."
<<
endl
;
}
}
}
template
<
class
T
>
void
check_mref_results
(
const
string
&
description
,
const
vector
<
T
>&
expected
,
mem_ref
<
T
>&
result
)
{
auto
exp_res
=
result
.
data
();
CAF_REQUIRE
(
exp_res
);
auto
res
=
*
exp_res
;
auto
cond
=
(
expected
==
res
);
CAF_CHECK
(
cond
);
if
(
!
cond
)
{
CAF_ERROR
(
description
<<
" failed."
);
cout
<<
"Expected: "
<<
endl
;
for
(
size_t
i
=
0
;
i
<
expected
.
size
();
++
i
)
{
cout
<<
" "
<<
expected
[
i
];
}
cout
<<
endl
<<
"Received: "
<<
endl
;
for
(
size_t
i
=
0
;
i
<
res
.
size
();
++
i
)
{
cout
<<
" "
<<
res
[
i
];
}
cout
<<
endl
;
}
}
void
test_opencl
(
actor_system
&
sys
)
{
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device_if
([](
const
device_ptr
){
return
true
;
});
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
const
ivec
expected1
{
56
,
62
,
68
,
74
,
152
,
174
,
196
,
218
,
248
,
286
,
324
,
362
,
344
,
398
,
452
,
506
};
auto
w1
=
mngr
.
spawn
(
prog
,
kn_matrix
,
opencl
::
nd_range
{
dims
{
matrix_size
,
matrix_size
}},
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{});
self
->
send
(
w1
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Simple matrix multiplication using vectors"
" (kernel wrapped in program)"
,
expected1
,
result
);
},
others
>>
wrong_msg
);
opencl
::
nd_range
range2
{
dims
{
matrix_size
,
matrix_size
}};
// Pass kernel directly to the actor
auto
w2
=
mngr
.
spawn
(
kernel_source
,
kn_matrix
,
range2
,
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{});
self
->
send
(
w2
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Simple matrix multiplication using vectors"
" (kernel passed directly)"
,
expected1
,
result
);
},
others
>>
wrong_msg
);
const
matrix_type
expected2
(
move
(
expected1
));
auto
map_arg
=
[](
message
&
msg
)
->
optional
<
message
>
{
return
msg
.
apply
(
[](
matrix_type
&
mx
)
{
return
make_message
(
move
(
mx
.
data
()));
}
);
};
auto
map_res
=
[](
ivec
result
)
->
message
{
return
make_message
(
matrix_type
{
move
(
result
)});
};
opencl
::
nd_range
range3
{
dims
{
matrix_size
,
matrix_size
}};
// let the runtime choose the device
auto
w3
=
mngr
.
spawn
(
mngr
.
create_program
(
kernel_source
),
kn_matrix
,
range3
,
map_arg
,
map_res
,
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{});
self
->
send
(
w3
,
make_iota_matrix
<
matrix_size
>
());
self
->
receive
(
[
&
](
const
matrix_type
&
result
)
{
check_vector_results
(
"Matrix multiplication with user defined type "
"(kernel wrapped in program)"
,
expected2
.
data
(),
result
.
data
());
},
others
>>
wrong_msg
);
opencl
::
nd_range
range4
{
dims
{
matrix_size
,
matrix_size
}};
auto
w4
=
mngr
.
spawn
(
prog
,
kn_matrix
,
range4
,
map_arg
,
map_res
,
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{});
self
->
send
(
w4
,
make_iota_matrix
<
matrix_size
>
());
self
->
receive
(
[
&
](
const
matrix_type
&
result
)
{
check_vector_results
(
"Matrix multiplication with user defined type"
,
expected2
.
data
(),
result
.
data
());
},
others
>>
wrong_msg
);
#ifndef CAF_NO_EXCEPTIONS
CAF_MESSAGE
(
"Expecting exception (compiling invalid kernel, "
"semicolon is missing)."
);
try
{
/* auto expected_error = */
mngr
.
create_program
(
kernel_source_error
);
}
catch
(
const
exception
&
exc
)
{
CAF_MESSAGE
(
"got: "
<<
exc
.
what
());
}
#endif // CAF_NO_EXCEPTIONS
// create program with opencl compiler flags
auto
prog5
=
mngr
.
create_program
(
kernel_source_compiler_flag
,
compiler_flag
);
opencl
::
nd_range
range5
{
dims
{
array_size
}};
auto
w5
=
mngr
.
spawn
(
prog5
,
kn_compiler_flag
,
range5
,
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{});
self
->
send
(
w5
,
make_iota_vector
<
int
>
(
array_size
));
auto
expected3
=
make_iota_vector
<
int
>
(
array_size
);
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Passing compiler flags"
,
expected3
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
max_work_item_sizes
()[
0
],
size_t
{
512
});
auto
reduce_buffer_size
=
static_cast
<
size_t
>
(
max_wg_size
)
*
8
;
auto
reduce_local_size
=
static_cast
<
size_t
>
(
max_wg_size
);
auto
reduce_work_groups
=
reduce_buffer_size
/
reduce_local_size
;
auto
reduce_global_size
=
reduce_buffer_size
;
auto
reduce_result_size
=
reduce_work_groups
;
ivec
arr6
(
reduce_buffer_size
);
int
n
=
static_cast
<
int
>
(
arr6
.
capacity
());
generate
(
arr6
.
begin
(),
arr6
.
end
(),
[
&
]{
return
--
n
;
});
opencl
::
nd_range
range6
{
dims
{
reduce_global_size
},
dims
{
},
// no offset
dims
{
reduce_local_size
}};
auto
result_size_6
=
[
reduce_result_size
](
const
ivec
&
)
{
return
reduce_result_size
;
};
auto
w6
=
mngr
.
spawn
(
prog
,
kn_reduce
,
range6
,
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{
result_size_6
});
self
->
send
(
w6
,
move
(
arr6
));
auto
wg_size_as_int
=
static_cast
<
int
>
(
max_wg_size
);
ivec
expected4
{
wg_size_as_int
*
7
,
wg_size_as_int
*
6
,
wg_size_as_int
*
5
,
wg_size_as_int
*
4
,
wg_size_as_int
*
3
,
wg_size_as_int
*
2
,
wg_size_as_int
,
0
};
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Passing size for the output"
,
expected4
,
result
);
},
others
>>
wrong_msg
);
// calculator function for getting the size of the output
auto
result_size_7
=
[](
const
ivec
&
)
{
return
problem_size
;
};
// constant memory arguments
const
ivec
arr7
{
static_cast
<
int
>
(
problem_size
)};
auto
w7
=
mngr
.
spawn
(
kernel_source
,
kn_const
,
opencl
::
nd_range
{
dims
{
problem_size
}},
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{
result_size_7
});
self
->
send
(
w7
,
move
(
arr7
));
ivec
expected5
(
problem_size
);
fill
(
begin
(
expected5
),
end
(
expected5
),
static_cast
<
int
>
(
problem_size
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Using const input argument"
,
expected5
,
result
);
},
others
>>
wrong_msg
);
}
void
test_arguments
(
actor_system
&
sys
)
{
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device_if
([](
const
device_ptr
){
return
true
;
});
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
const
ivec
expected1
{
56
,
62
,
68
,
74
,
152
,
174
,
196
,
218
,
248
,
286
,
324
,
362
,
344
,
398
,
452
,
506
};
auto
w1
=
mngr
.
spawn
(
mngr
.
create_program
(
kernel_source
,
""
,
dev
),
kn_matrix
,
opencl
::
nd_range
{
dims
{
matrix_size
,
matrix_size
}},
opencl
::
in
<
int
>
{},
opencl
::
out
<
int
>
{});
self
->
send
(
w1
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"arguments: from in to out"
,
expected1
,
result
);
},
others
>>
wrong_msg
);
ivec
input9
=
make_iota_vector
<
int
>
(
problem_size
);
ivec
expected9
{
input9
};
for_each
(
begin
(
expected9
),
end
(
expected9
),
[](
int
&
val
){
val
*=
2
;
});
auto
w9
=
mngr
.
spawn
(
kernel_source
,
kn_inout
,
nd_range
{
dims
{
problem_size
}},
opencl
::
in_out
<
int
>
{});
self
->
send
(
w9
,
move
(
input9
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing in_out arugment"
,
expected9
,
result
);
},
others
>>
wrong_msg
);
ivec
input10
=
make_iota_vector
<
int
>
(
problem_size
);
ivec
expected10
{
input10
};
for_each
(
begin
(
expected10
),
end
(
expected10
),
[](
int
&
val
){
val
*=
2
;
});
auto
result_size_10
=
[
=
](
const
ivec
&
input
)
{
return
input
.
size
();
};
auto
w10
=
mngr
.
spawn
(
kernel_source
,
kn_scratch
,
nd_range
{
dims
{
problem_size
}},
opencl
::
in_out
<
int
>
{},
opencl
::
scratch
<
int
>
{
result_size_10
});
self
->
send
(
w10
,
move
(
input10
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing buffer arugment"
,
expected10
,
result
);
},
others
>>
wrong_msg
);
// test local
size_t
la_global
=
256
;
size_t
la_local
=
la_global
/
2
;
ivec
input_local
=
make_iota_vector
<
int
>
(
la_global
);
ivec
expected_local
{
input_local
};
auto
last
=
0
;
for
(
size_t
i
=
0
;
i
<
la_global
;
++
i
)
{
if
(
i
==
la_local
)
last
=
0
;
auto
tmp
=
expected_local
[
i
];
expected_local
[
i
]
=
last
;
last
+=
tmp
;
}
auto
work_local
=
mngr
.
spawn
(
kernel_source
,
kn_local
,
nd_range
{
dims
{
la_global
},
{},
dims
{
la_local
}},
opencl
::
in_out
<
int
>
{},
opencl
::
local
<
int
>
{
la_local
});
self
->
send
(
work_local
,
std
::
move
(
input_local
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing local arugment"
,
expected_local
,
result
);
}
);
// Same test, different argument order
input_local
=
make_iota_vector
<
int
>
(
la_global
);
work_local
=
mngr
.
spawn
(
kernel_source
,
kn_order
,
nd_range
{
dims
{
la_global
},
{},
dims
{
la_local
}},
opencl
::
local
<
int
>
{
la_local
},
opencl
::
in_out
<
int
>
{});
self
->
send
(
work_local
,
std
::
move
(
input_local
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing local arugment"
,
expected_local
,
result
);
}
);
// Test private argument
ivec
input_private
=
make_iota_vector
<
int
>
(
problem_size
);
int
val_private
=
42
;
ivec
expected_private
{
input_private
};
for_each
(
begin
(
expected_private
),
end
(
expected_private
),
[
val_private
](
int
&
val
){
val
+=
val_private
;
});
auto
worker_private
=
mngr
.
spawn
(
kernel_source
,
kn_private
,
nd_range
{
dims
{
problem_size
}},
opencl
::
in_out
<
int
>
{},
opencl
::
priv
<
int
>
{
val_private
});
self
->
send
(
worker_private
,
std
::
move
(
input_private
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing private arugment"
,
expected_private
,
result
);
}
);
}
CAF_TEST
(
opencl_basics
)
{
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
()
.
add_message_type
<
ivec
>
(
"int_vector"
)
.
add_message_type
<
matrix_type
>
(
"square_matrix"
);
actor_system
system
{
cfg
};
test_opencl
(
system
);
system
.
await_all_actors_done
();
}
CAF_TEST
(
opencl_arguments
)
{
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
()
.
add_message_type
<
ivec
>
(
"int_vector"
)
.
add_message_type
<
matrix_type
>
(
"square_matrix"
);
actor_system
system
{
cfg
};
test_arguments
(
system
);
system
.
await_all_actors_done
();
}
CAF_TEST
(
opencl_mem_refs
)
{
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
();
actor_system
system
{
cfg
};
auto
&
mngr
=
system
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
// global arguments
vector
<
uint32_t
>
input
{
1
,
2
,
3
,
4
};
auto
buf_1
=
dev
->
global_argument
(
input
,
buffer_type
::
input_output
);
CAF_CHECK_EQUAL
(
buf_1
.
size
(),
input
.
size
());
auto
res_1
=
buf_1
.
data
();
CAF_CHECK
(
res_1
);
CAF_CHECK_EQUAL
(
res_1
->
size
(),
input
.
size
());
check_vector_results
(
"Testing mem_ref"
,
input
,
*
res_1
);
auto
res_2
=
buf_1
.
data
(
2ul
);
CAF_CHECK
(
res_2
);
CAF_CHECK_EQUAL
(
res_2
->
size
(),
2ul
);
CAF_CHECK_EQUAL
((
*
res_2
)[
0
],
input
[
0
]);
CAF_CHECK_EQUAL
((
*
res_2
)[
1
],
input
[
1
]);
vector
<
uint32_t
>
new_input
{
1
,
2
,
3
,
4
,
5
};
buf_1
=
dev
->
global_argument
(
new_input
,
buffer_type
::
input_output
);
CAF_CHECK_EQUAL
(
buf_1
.
size
(),
new_input
.
size
());
auto
res_3
=
buf_1
.
data
();
CAF_CHECK
(
res_3
);
mem_ref
<
uint32_t
>
buf_2
{
std
::
move
(
buf_1
)};
CAF_CHECK_EQUAL
(
buf_2
.
size
(),
new_input
.
size
());
auto
res_4
=
buf_2
.
data
();
CAF_CHECK
(
res_4
);
buf_2
.
reset
();
auto
res_5
=
buf_2
.
data
();
CAF_CHECK
(
!
res_5
);
}
CAF_TEST
(
opencl_argument_info
)
{
using
base_t
=
int
;
using
in_arg_t
=
::
type_list
<
opencl
::
in
<
base_t
>>
;
using
in_arg_info_t
=
typename
cl_arg_info_list
<
in_arg_t
>::
type
;
using
in_arg_wrap_t
=
typename
::
tl_head
<
in_arg_info_t
>::
type
;
static_assert
(
in_arg_wrap_t
::
in_pos
==
0
,
"In-index for `in` wrong."
);
static_assert
(
in_arg_wrap_t
::
out_pos
==
-
1
,
"Out-index for `in` wrong."
);
using
out_arg_t
=
::
type_list
<
opencl
::
out
<
base_t
>>
;
using
out_arg_info_t
=
typename
cl_arg_info_list
<
out_arg_t
>::
type
;
using
out_arg_wrap_t
=
typename
::
tl_head
<
out_arg_info_t
>::
type
;
static_assert
(
out_arg_wrap_t
::
in_pos
==
-
1
,
"In-index for `out` wrong."
);
static_assert
(
out_arg_wrap_t
::
out_pos
==
0
,
"Out-index for `out` wrong."
);
using
io_arg_t
=
::
type_list
<
opencl
::
in_out
<
base_t
>>
;
using
io_arg_info_t
=
typename
cl_arg_info_list
<
io_arg_t
>::
type
;
using
io_arg_wrap_t
=
typename
::
tl_head
<
io_arg_info_t
>::
type
;
static_assert
(
io_arg_wrap_t
::
in_pos
==
0
,
"In-index for `in_out` wrong."
);
static_assert
(
io_arg_wrap_t
::
out_pos
==
0
,
"Out-index for `in_out` wrong."
);
using
arg_list_t
=
::
type_list
<
opencl
::
in
<
base_t
>
,
opencl
::
out
<
base_t
>
,
opencl
::
local
<
base_t
>
,
opencl
::
in_out
<
base_t
>
,
opencl
::
priv
<
base_t
>
,
opencl
::
priv
<
base_t
,
val
>>
;
using
arg_info_list_t
=
typename
cl_arg_info_list
<
arg_list_t
>::
type
;
using
arg_info_0_t
=
typename
::
tl_at
<
arg_info_list_t
,
0
>::
type
;
static_assert
(
arg_info_0_t
::
in_pos
==
0
,
"In-index for `in` wrong."
);
static_assert
(
arg_info_0_t
::
out_pos
==
-
1
,
"Out-index for `in` wrong."
);
using
arg_info_1_t
=
typename
::
tl_at
<
arg_info_list_t
,
1
>::
type
;
static_assert
(
arg_info_1_t
::
in_pos
==
-
1
,
"In-index for `out` wrong."
);
static_assert
(
arg_info_1_t
::
out_pos
==
0
,
"Out-index for `out` wrong."
);
using
arg_info_2_t
=
typename
::
tl_at
<
arg_info_list_t
,
2
>::
type
;
static_assert
(
arg_info_2_t
::
in_pos
==
-
1
,
"In-index for `local` wrong."
);
static_assert
(
arg_info_2_t
::
out_pos
==
-
1
,
"Out-index for `local` wrong."
);
using
arg_info_3_t
=
typename
::
tl_at
<
arg_info_list_t
,
3
>::
type
;
static_assert
(
arg_info_3_t
::
in_pos
==
1
,
"In-index for `in_out` wrong."
);
static_assert
(
arg_info_3_t
::
out_pos
==
1
,
"Out-index for `in_out` wrong."
);
using
arg_info_4_t
=
typename
::
tl_at
<
arg_info_list_t
,
4
>::
type
;
static_assert
(
arg_info_4_t
::
in_pos
==
-
1
,
"In-index for `priv` wrong."
);
static_assert
(
arg_info_4_t
::
out_pos
==
-
1
,
"Out-index for `priv` wrong."
);
using
arg_info_5_t
=
typename
::
tl_at
<
arg_info_list_t
,
5
>::
type
;
static_assert
(
arg_info_5_t
::
in_pos
==
2
,
"In-index for `priv` wrong."
);
static_assert
(
arg_info_5_t
::
out_pos
==
-
1
,
"Out-index for `priv` wrong."
);
// gives the test some output.
CAF_CHECK_EQUAL
(
true
,
true
);
}
void
test_in_val_out_val
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing in: val -> out: val "
);
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
const
ivec
res1
{
56
,
62
,
68
,
74
,
152
,
174
,
196
,
218
,
248
,
286
,
324
,
362
,
344
,
398
,
452
,
506
};
auto
conf
=
opencl
::
nd_range
{
dims
{
matrix_size
,
matrix_size
}};
auto
w1
=
mngr
.
spawn
(
prog
,
kn_matrix
,
conf
,
in
<
int
>
{},
out
<
int
>
{});
self
->
send
(
w1
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Simple matrix multiplication using vectors"
" (kernel wrapped in program)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// Pass kernel directly to the actor
auto
w2
=
mngr
.
spawn
(
kernel_source
,
kn_matrix
,
conf
,
in
<
int
>
{},
out
<
int
>
{});
self
->
send
(
w2
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Simple matrix multiplication using vectors"
" (kernel passed directly)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// Wrap message in user-defined type and use mapping functions
const
matrix_type
res2
(
move
(
res1
));
auto
map_arg
=
[](
message
&
msg
)
->
optional
<
message
>
{
return
msg
.
apply
([](
matrix_type
&
mx
)
{
return
make_message
(
move
(
mx
.
data
()));
});
};
auto
map_res
=
[](
ivec
result
)
->
message
{
return
make_message
(
matrix_type
{
move
(
result
)});
};
auto
w3
=
mngr
.
spawn
(
prog
,
kn_matrix
,
conf
,
map_arg
,
map_res
,
in
<
int
,
val
>
{},
out
<
int
,
val
>
{});
self
->
send
(
w3
,
make_iota_matrix
<
matrix_size
>
());
self
->
receive
([
&
](
const
matrix_type
&
result
)
{
check_vector_results
(
"Matrix multiplication with user defined type "
"(kernel wrapped in program)"
,
res2
.
data
(),
result
.
data
());
},
others
>>
wrong_msg
);
// create program with opencl compiler flags
auto
prog2
=
mngr
.
create_program
(
kernel_source_compiler_flag
,
compiler_flag
);
nd_range
range2
{
dims
{
array_size
}};
auto
w4
=
mngr
.
spawn
(
prog2
,
kn_compiler_flag
,
range2
,
in
<
int
>
{},
out
<
int
>
{});
self
->
send
(
w4
,
make_iota_vector
<
int
>
(
array_size
));
auto
res3
=
make_iota_vector
<
int
>
(
array_size
);
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Passing compiler flags"
,
res3
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
max_work_item_sizes
()[
0
],
size_t
{
512
});
auto
reduce_buffer_size
=
static_cast
<
size_t
>
(
max_wg_size
)
*
8
;
auto
reduce_local_size
=
static_cast
<
size_t
>
(
max_wg_size
);
auto
reduce_work_groups
=
reduce_buffer_size
/
reduce_local_size
;
auto
reduce_global_size
=
reduce_buffer_size
;
auto
reduce_result_size
=
reduce_work_groups
;
ivec
input
(
reduce_buffer_size
);
int
n
=
static_cast
<
int
>
(
input
.
capacity
());
generate
(
input
.
begin
(),
input
.
end
(),
[
&
]{
return
--
n
;
});
nd_range
range3
{
dims
{
reduce_global_size
},
dims
{},
dims
{
reduce_local_size
}};
auto
res_size
=
[
&
](
const
ivec
&
)
{
return
reduce_result_size
;
};
auto
w5
=
mngr
.
spawn
(
prog
,
kn_reduce
,
range3
,
in
<
int
>
{},
out
<
int
>
{
res_size
});
self
->
send
(
w5
,
move
(
input
));
auto
wg_size_as_int
=
static_cast
<
int
>
(
max_wg_size
);
ivec
res4
{
wg_size_as_int
*
7
,
wg_size_as_int
*
6
,
wg_size_as_int
*
5
,
wg_size_as_int
*
4
,
wg_size_as_int
*
3
,
wg_size_as_int
*
2
,
wg_size_as_int
,
0
};
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Passing size for the output"
,
res4
,
result
);
},
others
>>
wrong_msg
);
// calculator function for getting the size of the output
auto
res_size2
=
[](
const
ivec
&
)
{
return
problem_size
;
};
// constant memory arguments
const
ivec
input2
{
static_cast
<
int
>
(
problem_size
)};
auto
w6
=
mngr
.
spawn
(
kernel_source
,
kn_const
,
nd_range
{
dims
{
problem_size
}},
in
<
int
>
{},
out
<
int
>
{
res_size2
});
self
->
send
(
w6
,
move
(
input2
));
ivec
res5
(
problem_size
);
fill
(
begin
(
res5
),
end
(
res5
),
static_cast
<
int
>
(
problem_size
));
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Using const input argument"
,
res5
,
result
);
},
others
>>
wrong_msg
);
}
void
test_in_val_out_mref
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing in: val -> out: mref"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
// tests
const
ivec
res1
{
56
,
62
,
68
,
74
,
152
,
174
,
196
,
218
,
248
,
286
,
324
,
362
,
344
,
398
,
452
,
506
};
auto
range
=
opencl
::
nd_range
{
dims
{
matrix_size
,
matrix_size
}};
auto
w1
=
mngr
.
spawn
(
prog
,
kn_matrix
,
range
,
in
<
int
>
{},
out
<
int
,
mref
>
{});
self
->
send
(
w1
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Simple matrix multiplication using vectors"
" (kernel wrapped in program)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// Pass kernel directly to the actor
auto
w2
=
mngr
.
spawn
(
kernel_source
,
kn_matrix
,
range
,
in
<
int
>
{},
out
<
int
,
mref
>
{});
self
->
send
(
w2
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Simple matrix multiplication using vectors"
" (kernel passed directly)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
max_work_item_sizes
()[
0
],
size_t
{
512
});
auto
reduce_buffer_size
=
static_cast
<
size_t
>
(
max_wg_size
)
*
8
;
auto
reduce_local_size
=
static_cast
<
size_t
>
(
max_wg_size
);
auto
reduce_work_groups
=
reduce_buffer_size
/
reduce_local_size
;
auto
reduce_global_size
=
reduce_buffer_size
;
auto
reduce_result_size
=
reduce_work_groups
;
ivec
input
(
reduce_buffer_size
);
int
n
=
static_cast
<
int
>
(
input
.
capacity
());
generate
(
input
.
begin
(),
input
.
end
(),
[
&
]{
return
--
n
;
});
nd_range
range3
{
dims
{
reduce_global_size
},
dims
{},
dims
{
reduce_local_size
}};
auto
res_size
=
[
&
](
const
ivec
&
)
{
return
reduce_result_size
;
};
auto
w5
=
mngr
.
spawn
(
prog
,
kn_reduce
,
range3
,
in
<
int
>
{},
out
<
int
,
mref
>
{
res_size
});
self
->
send
(
w5
,
move
(
input
));
auto
wg_size_as_int
=
static_cast
<
int
>
(
max_wg_size
);
ivec
res4
{
wg_size_as_int
*
7
,
wg_size_as_int
*
6
,
wg_size_as_int
*
5
,
wg_size_as_int
*
4
,
wg_size_as_int
*
3
,
wg_size_as_int
*
2
,
wg_size_as_int
*
1
,
wg_size_as_int
*
0
};
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Passing size for the output"
,
res4
,
result
);
},
others
>>
wrong_msg
);
}
void
test_in_mref_out_val
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing in: mref -> out: val "
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
// tests
const
ivec
res1
{
56
,
62
,
68
,
74
,
152
,
174
,
196
,
218
,
248
,
286
,
324
,
362
,
344
,
398
,
452
,
506
};
auto
range
=
opencl
::
nd_range
{
dims
{
matrix_size
,
matrix_size
}};
auto
w1
=
mngr
.
spawn
(
prog
,
kn_matrix
,
range
,
in
<
int
,
mref
>
{},
out
<
int
>
{});
auto
matrix1
=
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
);
auto
input1
=
dev
->
global_argument
(
matrix1
);
self
->
send
(
w1
,
input1
);
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Simple matrix multiplication using vectors"
" (kernel wrapped in program)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// Pass kernel directly to the actor
auto
w2
=
mngr
.
spawn
(
kernel_source
,
kn_matrix
,
range
,
in
<
int
,
mref
>
{},
out
<
int
,
val
>
{});
self
->
send
(
w2
,
input1
);
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Simple matrix multiplication using vectors"
" (kernel passed directly)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
max_work_item_sizes
()[
0
],
size_t
{
512
});
auto
reduce_buffer_size
=
static_cast
<
size_t
>
(
max_wg_size
)
*
8
;
auto
reduce_local_size
=
static_cast
<
size_t
>
(
max_wg_size
);
auto
reduce_work_groups
=
reduce_buffer_size
/
reduce_local_size
;
auto
reduce_global_size
=
reduce_buffer_size
;
auto
reduce_result_size
=
reduce_work_groups
;
ivec
values
(
reduce_buffer_size
);
int
n
=
static_cast
<
int
>
(
values
.
capacity
());
generate
(
values
.
begin
(),
values
.
end
(),
[
&
]{
return
--
n
;
});
nd_range
range3
{
dims
{
reduce_global_size
},
dims
{},
dims
{
reduce_local_size
}};
auto
res_size
=
[
&
](
const
iref
&
)
{
return
reduce_result_size
;
};
auto
w5
=
mngr
.
spawn
(
prog
,
kn_reduce
,
range3
,
in
<
int
,
mref
>
{},
out
<
int
>
{
res_size
});
auto
input2
=
dev
->
global_argument
(
values
);
self
->
send
(
w5
,
input2
);
auto
multiplier
=
static_cast
<
int
>
(
max_wg_size
);
ivec
res4
{
multiplier
*
7
,
multiplier
*
6
,
multiplier
*
5
,
multiplier
*
4
,
multiplier
*
3
,
multiplier
*
2
,
multiplier
*
1
,
multiplier
*
0
};
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Passing size for the output"
,
res4
,
result
);
},
others
>>
wrong_msg
);
}
void
test_in_mref_out_mref
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing in: mref -> out: mref"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
// tests
const
ivec
res1
{
56
,
62
,
68
,
74
,
152
,
174
,
196
,
218
,
248
,
286
,
324
,
362
,
344
,
398
,
452
,
506
};
auto
range
=
opencl
::
nd_range
{
dims
{
matrix_size
,
matrix_size
}};
auto
w1
=
mngr
.
spawn
(
prog
,
kn_matrix
,
range
,
in
<
int
,
mref
>
{},
out
<
int
,
mref
>
{});
auto
matrix1
=
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
);
auto
input1
=
dev
->
global_argument
(
matrix1
);
self
->
send
(
w1
,
input1
);
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Simple matrix multiplication using vectors"
" (kernel wrapped in program)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// Pass kernel directly to the actor
auto
w2
=
mngr
.
spawn
(
kernel_source
,
kn_matrix
,
range
,
in
<
int
,
mref
>
{},
out
<
int
,
mref
>
{});
self
->
send
(
w2
,
input1
);
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Simple matrix multiplication using vectors"
" (kernel passed directly)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
max_work_item_sizes
()[
0
],
size_t
{
512
});
auto
reduce_buffer_size
=
static_cast
<
size_t
>
(
max_wg_size
)
*
8
;
auto
reduce_local_size
=
static_cast
<
size_t
>
(
max_wg_size
);
auto
reduce_work_groups
=
reduce_buffer_size
/
reduce_local_size
;
auto
reduce_global_size
=
reduce_buffer_size
;
auto
reduce_result_size
=
reduce_work_groups
;
ivec
values
(
reduce_buffer_size
);
int
n
=
static_cast
<
int
>
(
values
.
capacity
());
generate
(
values
.
begin
(),
values
.
end
(),
[
&
]{
return
--
n
;
});
nd_range
range3
{
dims
{
reduce_global_size
},
dims
{},
dims
{
reduce_local_size
}};
auto
res_size
=
[
&
](
const
iref
&
)
{
return
reduce_result_size
;
};
auto
w5
=
mngr
.
spawn
(
prog
,
kn_reduce
,
range3
,
in
<
int
,
mref
>
{},
out
<
int
,
mref
>
{
res_size
});
auto
input2
=
dev
->
global_argument
(
values
);
self
->
send
(
w5
,
input2
);
auto
multiplier
=
static_cast
<
int
>
(
max_wg_size
);
ivec
res4
{
multiplier
*
7
,
multiplier
*
6
,
multiplier
*
5
,
multiplier
*
4
,
multiplier
*
3
,
multiplier
*
2
,
multiplier
,
0
};
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Passing size for the output"
,
res4
,
result
);
},
others
>>
wrong_msg
);
}
void
test_varying_arguments
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing varying argument order "
"(Might fail on some integrated GPUs)"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
// tests
size_t
size
=
23
;
nd_range
range
{
dims
{
size
}};
auto
input1
=
make_iota_vector
<
int
>
(
size
);
auto
input2
=
dev
->
global_argument
(
input1
);
auto
w1
=
mngr
.
spawn
(
prog
,
kn_varying
,
range
,
in
<
int
>
{},
out
<
int
>
{},
in
<
int
>
{},
out
<
int
>
{});
self
->
send
(
w1
,
input1
,
input1
);
self
->
receive
([
&
](
const
ivec
&
res1
,
const
ivec
&
res2
)
{
check_vector_results
(
"Varying args (vec only), output 1"
,
input1
,
res1
);
check_vector_results
(
"Varying args (vec only), output 2"
,
input1
,
res2
);
},
others
>>
wrong_msg
);
auto
w2
=
mngr
.
spawn
(
prog
,
kn_varying
,
range
,
in
<
int
,
mref
>
{},
out
<
int
>
{},
in
<
int
>
{},
out
<
int
,
mref
>
{});
self
->
send
(
w2
,
input2
,
input1
);
self
->
receive
([
&
](
const
ivec
&
res1
,
iref
&
res2
)
{
check_vector_results
(
"Varying args (vec), output 1"
,
input1
,
res1
);
check_mref_results
(
"Varying args (ref), output 2"
,
input1
,
res2
);
},
others
>>
wrong_msg
);
}
void
test_inout
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing in_out arguments"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
// tests
ivec
input
=
make_iota_vector
<
int
>
(
problem_size
);
auto
input2
=
dev
->
global_argument
(
input
);
auto
input3
=
dev
->
global_argument
(
input
);
ivec
res
{
input
};
for_each
(
begin
(
res
),
end
(
res
),
[](
int
&
val
){
val
*=
2
;
});
auto
range
=
nd_range
{
dims
{
problem_size
}};
auto
w1
=
mngr
.
spawn
(
kernel_source
,
kn_inout
,
range
,
in_out
<
int
,
val
,
val
>
{});
self
->
send
(
w1
,
input
);
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing in_out (val -> val)"
,
res
,
result
);
},
others
>>
wrong_msg
);
auto
w2
=
mngr
.
spawn
(
kernel_source
,
kn_inout
,
range
,
in_out
<
int
,
val
,
mref
>
{});
self
->
send
(
w2
,
input
);
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Testing in_out (val -> mref)"
,
res
,
result
);
},
others
>>
wrong_msg
);
auto
w3
=
mngr
.
spawn
(
kernel_source
,
kn_inout
,
range
,
in_out
<
int
,
mref
,
val
>
{});
self
->
send
(
w3
,
input2
);
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing in_out (mref -> val)"
,
res
,
result
);
},
others
>>
wrong_msg
);
auto
w4
=
mngr
.
spawn
(
kernel_source
,
kn_inout
,
range
,
in_out
<
int
,
mref
,
mref
>
{});
self
->
send
(
w4
,
input3
);
self
->
receive
([
&
](
iref
&
result
)
{
check_mref_results
(
"Testing in_out (mref -> mref)"
,
res
,
result
);
},
others
>>
wrong_msg
);
}
void
test_priv
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing priv argument"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
// tests
nd_range
range
{
dims
{
problem_size
}};
ivec
input
=
make_iota_vector
<
int
>
(
problem_size
);
int
value
=
42
;
ivec
res
{
input
};
for_each
(
begin
(
res
),
end
(
res
),
[
&
](
int
&
val
){
val
+=
value
;
});
auto
w1
=
mngr
.
spawn
(
kernel_source
,
kn_private
,
range
,
in_out
<
int
>
{},
priv
<
int
>
{
value
});
self
->
send
(
w1
,
input
);
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing hidden private arugment"
,
res
,
result
);
},
others
>>
wrong_msg
);
auto
w2
=
mngr
.
spawn
(
kernel_source
,
kn_private
,
range
,
in_out
<
int
>
{},
priv
<
int
,
val
>
{});
self
->
send
(
w2
,
input
,
value
);
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing val private arugment"
,
res
,
result
);
},
others
>>
wrong_msg
);
}
void
test_local
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing local argument"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
find_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
scoped_actor
self
{
sys
};
auto
wrong_msg
=
[
&
](
message_view
&
x
)
->
result
<
message
>
{
CAF_ERROR
(
"unexpected message"
<<
x
.
content
().
stringify
());
return
sec
::
unexpected_message
;
};
// tests
size_t
global_size
=
256
;
size_t
local_size
=
global_size
/
2
;
ivec
res
=
make_iota_vector
<
int
>
(
global_size
);
auto
last
=
0
;
for
(
size_t
i
=
0
;
i
<
global_size
;
++
i
)
{
if
(
i
==
local_size
)
last
=
0
;
auto
tmp
=
res
[
i
];
res
[
i
]
=
last
;
last
+=
tmp
;
}
auto
range
=
nd_range
{
dims
{
global_size
},
{},
dims
{
local_size
}};
auto
w
=
mngr
.
spawn
(
kernel_source
,
kn_local
,
range
,
in_out
<
int
>
{},
local
<
int
>
{
local_size
});
self
->
send
(
w
,
make_iota_vector
<
int
>
(
global_size
));
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing local arugment"
,
res
,
result
);
},
others
>>
wrong_msg
);
// Same test, different argument order
w
=
mngr
.
spawn
(
kernel_source
,
kn_order
,
range
,
local
<
int
>
{
local_size
},
in_out
<
int
>
{});
self
->
send
(
w
,
make_iota_vector
<
int
>
(
global_size
));
self
->
receive
([
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Testing local arugment"
,
res
,
result
);
},
others
>>
wrong_msg
);
}
CAF_TEST
(
actor_facade
)
{
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
()
.
add_message_type
<
ivec
>
(
"int_vector"
)
.
add_message_type
<
matrix_type
>
(
"square_matrix"
);
actor_system
system
{
cfg
};
test_in_val_out_val
(
system
);
test_in_val_out_mref
(
system
);
test_in_mref_out_val
(
system
);
test_in_mref_out_mref
(
system
);
test_varying_arguments
(
system
);
test_inout
(
system
);
test_priv
(
system
);
test_local
(
system
);
system
.
await_all_actors_done
();
}
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment