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
108e51ff
Commit
108e51ff
authored
May 27, 2016
by
Joseph Noir
Browse files
Options
Browse Files
Download
Plain Diff
Adapt to actor_system updates
parents
785153d3
36b3f4f3
Changes
21
Expand all
Show whitespace changes
Inline
Side-by-side
Showing
21 changed files
with
635 additions
and
863 deletions
+635
-863
libcaf_opencl/CMakeLists.txt
libcaf_opencl/CMakeLists.txt
+2
-2
libcaf_opencl/caf/opencl/actor_facade.hpp
libcaf_opencl/caf/opencl/actor_facade.hpp
+60
-35
libcaf_opencl/caf/opencl/all.hpp
libcaf_opencl/caf/opencl/all.hpp
+1
-1
libcaf_opencl/caf/opencl/arguments.hpp
libcaf_opencl/caf/opencl/arguments.hpp
+5
-5
libcaf_opencl/caf/opencl/command.hpp
libcaf_opencl/caf/opencl/command.hpp
+30
-21
libcaf_opencl/caf/opencl/detail/spawn_helper.hpp
libcaf_opencl/caf/opencl/detail/spawn_helper.hpp
+61
-0
libcaf_opencl/caf/opencl/device.hpp
libcaf_opencl/caf/opencl/device.hpp
+3
-1
libcaf_opencl/caf/opencl/manager.hpp
libcaf_opencl/caf/opencl/manager.hpp
+179
-0
libcaf_opencl/caf/opencl/opencl_err.hpp
libcaf_opencl/caf/opencl/opencl_err.hpp
+1
-1
libcaf_opencl/caf/opencl/platform.hpp
libcaf_opencl/caf/opencl/platform.hpp
+1
-2
libcaf_opencl/caf/opencl/program.hpp
libcaf_opencl/caf/opencl/program.hpp
+5
-15
libcaf_opencl/caf/opencl/spawn_cl.hpp
libcaf_opencl/caf/opencl/spawn_cl.hpp
+0
-322
libcaf_opencl/examples/proper_matrix.cpp
libcaf_opencl/examples/proper_matrix.cpp
+25
-16
libcaf_opencl/examples/simple_matrix.cpp
libcaf_opencl/examples/simple_matrix.cpp
+14
-10
libcaf_opencl/src/device.cpp
libcaf_opencl/src/device.cpp
+2
-2
libcaf_opencl/src/manager.cpp
libcaf_opencl/src/manager.cpp
+175
-0
libcaf_opencl/src/opencl_err.cpp
libcaf_opencl/src/opencl_err.cpp
+4
-3
libcaf_opencl/src/platform.cpp
libcaf_opencl/src/platform.cpp
+1
-1
libcaf_opencl/src/program.cpp
libcaf_opencl/src/program.cpp
+6
-56
libcaf_opencl/test/opencl.cpp
libcaf_opencl/test/opencl.cpp
+60
-78
libcaf_opencl/test/opencl_deprecated.cpp
libcaf_opencl/test/opencl_deprecated.cpp
+0
-292
No files found.
libcaf_opencl/CMakeLists.txt
View file @
108e51ff
...
@@ -4,13 +4,14 @@ project(caf_opencl C CXX)
...
@@ -4,13 +4,14 @@ project(caf_opencl C CXX)
# get header files; only needed by CMake generators,
# get header files; only needed by CMake generators,
# e.g., for creating proper Xcode projects
# e.g., for creating proper Xcode projects
file
(
GLOB LIBCAF_OPENCL_HDRS
"caf/opencl/*.hpp"
)
file
(
GLOB LIBCAF_OPENCL_HDRS
"caf/opencl/*.hpp"
)
file
(
GLOB LIBCAF_OPENCL_HDRS
"caf/opencl/detail/*.hpp"
)
add_custom_target
(
libcaf_opencl
)
add_custom_target
(
libcaf_opencl
)
# list cpp files excluding platform-dependent files
# list cpp files excluding platform-dependent files
set
(
LIBCAF_OPENCL_SRCS
set
(
LIBCAF_OPENCL_SRCS
src/global.cpp
src/global.cpp
src/m
etainfo
.cpp
src/m
anager
.cpp
src/program.cpp
src/program.cpp
src/opencl_err.cpp
src/opencl_err.cpp
src/platform.cpp
src/platform.cpp
...
@@ -44,4 +45,3 @@ link_directories(${LD_DIRS})
...
@@ -44,4 +45,3 @@ link_directories(${LD_DIRS})
include_directories
(
.
${
INCLUDE_DIRS
}
)
include_directories
(
.
${
INCLUDE_DIRS
}
)
# install includes
# install includes
install
(
DIRECTORY caf/ DESTINATION include/caf FILES_MATCHING PATTERN
"*.hpp"
)
install
(
DIRECTORY caf/ DESTINATION include/caf FILES_MATCHING PATTERN
"*.hpp"
)
install
(
DIRECTORY cppa/ DESTINATION include/cppa FILES_MATCHING PATTERN
"*.hpp"
)
libcaf_opencl/caf/opencl/actor_facade.hpp
View file @
108e51ff
...
@@ -27,7 +27,6 @@
...
@@ -27,7 +27,6 @@
#include "caf/all.hpp"
#include "caf/all.hpp"
#include "caf/channel.hpp"
#include "caf/intrusive_ptr.hpp"
#include "caf/intrusive_ptr.hpp"
#include "caf/detail/int_list.hpp"
#include "caf/detail/int_list.hpp"
...
@@ -45,7 +44,7 @@
...
@@ -45,7 +44,7 @@
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
class
m
etainfo
;
class
m
anager
;
template
<
class
List
>
template
<
class
List
>
struct
function_sig_from_outputs
;
struct
function_sig_from_outputs
;
...
@@ -64,7 +63,7 @@ struct command_sig_from_outputs<T, detail::type_list<Ts...>> {
...
@@ -64,7 +63,7 @@ struct command_sig_from_outputs<T, detail::type_list<Ts...>> {
};
};
template
<
class
...
Ts
>
template
<
class
...
Ts
>
class
actor_facade
:
public
abstract
_actor
{
class
actor_facade
:
public
monitorable
_actor
{
public:
public:
using
arg_types
=
detail
::
type_list
<
Ts
...
>
;
using
arg_types
=
detail
::
type_list
<
Ts
...
>
;
using
unpacked_types
=
typename
detail
::
tl_map
<
arg_types
,
extract_type
>::
type
;
using
unpacked_types
=
typename
detail
::
tl_map
<
arg_types
,
extract_type
>::
type
;
...
@@ -73,7 +72,7 @@ public:
...
@@ -73,7 +72,7 @@ public:
typename
detail
::
tl_filter
<
arg_types
,
is_input_arg
>::
type
;
typename
detail
::
tl_filter
<
arg_types
,
is_input_arg
>::
type
;
using
input_types
=
using
input_types
=
typename
detail
::
tl_map
<
input_wrapped_types
,
extract_type
>::
type
;
typename
detail
::
tl_map
<
input_wrapped_types
,
extract_type
>::
type
;
using
input_mapping
=
std
::
function
<
maybe
<
message
>
(
message
&
)
>
;
using
input_mapping
=
std
::
function
<
optional
<
message
>
(
message
&
)
>
;
using
output_wrapped_types
=
using
output_wrapped_types
=
typename
detail
::
tl_filter
<
arg_types
,
is_output_arg
>::
type
;
typename
detail
::
tl_filter
<
arg_types
,
is_output_arg
>::
type
;
...
@@ -90,40 +89,57 @@ public:
...
@@ -90,40 +89,57 @@ public:
using
command_type
=
using
command_type
=
typename
command_sig_from_outputs
<
actor_facade
,
output_types
>::
type
;
typename
command_sig_from_outputs
<
actor_facade
,
output_types
>::
type
;
static
intrusive_ptr
<
actor_facade
>
create
(
const
program
&
prog
,
const
char
*
name
()
const
override
{
const
char
*
kernel_name
,
return
"OpenCL actor"
;
const
spawn_config
&
config
,
}
input_mapping
map_args
,
output_mapping
map_result
,
static
actor
create
(
actor_config
actor_cfg
,
const
program
&
prog
,
const
char
*
kernel_name
,
const
spawn_config
&
spawn_cfg
,
input_mapping
map_args
,
output_mapping
map_result
,
Ts
&&
...
xs
)
{
Ts
&&
...
xs
)
{
if
(
confi
g
.
dimensions
().
empty
())
{
if
(
spawn_cf
g
.
dimensions
().
empty
())
{
auto
str
=
"OpenCL kernel needs at least 1 global dimension."
;
auto
str
=
"OpenCL kernel needs at least 1 global dimension."
;
CAF_LOG
F
_ERROR
(
str
);
CAF_LOG_ERROR
(
str
);
throw
std
::
runtime_error
(
str
);
throw
std
::
runtime_error
(
str
);
}
}
auto
check_vec
=
[
&
](
const
dim_vec
&
vec
,
const
char
*
name
)
{
auto
check_vec
=
[
&
](
const
dim_vec
&
vec
,
const
char
*
name
)
{
if
(
!
vec
.
empty
()
&&
vec
.
size
()
!=
confi
g
.
dimensions
().
size
())
{
if
(
!
vec
.
empty
()
&&
vec
.
size
()
!=
spawn_cf
g
.
dimensions
().
size
())
{
std
::
ostringstream
oss
;
std
::
ostringstream
oss
;
oss
<<
name
<<
" vector is not empty, but "
oss
<<
name
<<
" vector is not empty, but "
<<
"its size differs from global dimensions vector's size"
;
<<
"its size differs from global dimensions vector's size"
;
CAF_LOG
F_ERROR
(
oss
.
str
(
));
CAF_LOG
_ERROR
(
CAF_ARG
(
oss
.
str
()
));
throw
std
::
runtime_error
(
oss
.
str
());
throw
std
::
runtime_error
(
oss
.
str
());
}
}
};
};
check_vec
(
config
.
offsets
(),
"offsets"
);
check_vec
(
spawn_cfg
.
offsets
(),
"offsets"
);
check_vec
(
config
.
local_dimensions
(),
"local dimensions"
);
check_vec
(
spawn_cfg
.
local_dimensions
(),
"local dimensions"
);
cl_int
err
=
0
;
auto
&
sys
=
actor_cfg
.
host
->
system
();
auto
itr
=
prog
.
available_kernels_
.
find
(
kernel_name
);
if
(
itr
==
prog
.
available_kernels_
.
end
())
{
kernel_ptr
kernel
;
kernel_ptr
kernel
;
kernel
.
reset
(
clCreateKernel
(
prog
.
program_
.
get
(),
kernel_name
,
&
err
),
false
);
kernel
.
reset
(
v2get
(
CAF_CLF
(
clCreateKernel
),
prog
.
program_
.
get
(),
if
(
err
!=
CL_SUCCESS
)
kernel_name
),
return
nullptr
;
false
);
return
new
actor_facade
(
prog
,
kernel
,
config
,
return
make_actor
<
actor_facade
,
actor
>
(
sys
.
next_actor_id
(),
sys
.
node
(),
std
::
move
(
map_args
),
std
::
move
(
map_result
),
&
sys
,
std
::
move
(
actor_cfg
),
prog
,
kernel
,
spawn_cfg
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward_as_tuple
(
xs
...));
std
::
forward_as_tuple
(
xs
...));
}
else
{
return
make_actor
<
actor_facade
,
actor
>
(
sys
.
next_actor_id
(),
sys
.
node
(),
&
sys
,
std
::
move
(
actor_cfg
),
prog
,
itr
->
second
,
spawn_cfg
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward_as_tuple
(
xs
...));
}
}
}
void
enqueue
(
const
actor_addr
&
sender
,
message_id
mid
,
message
content
,
void
enqueue
(
strong_actor_ptr
sender
,
message_id
mid
,
message
content
,
execution_unit
*
)
override
{
execution_unit
*
)
override
{
CAF_PUSH_AID
(
id
());
CAF_LOG_TRACE
(
""
);
CAF_LOG_TRACE
(
""
);
if
(
map_args_
)
{
if
(
map_args_
)
{
auto
mapped
=
map_args_
(
content
);
auto
mapped
=
map_args_
(
content
);
...
@@ -135,14 +151,15 @@ public:
...
@@ -135,14 +151,15 @@ public:
if
(
!
content
.
match_elements
(
input_types
{}))
{
if
(
!
content
.
match_elements
(
input_types
{}))
{
return
;
return
;
}
}
response_promise
hdl
{
this
->
address
(),
sender
,
mid
.
response_id
()}
;
auto
hdl
=
std
::
make_tuple
(
sender
,
mid
.
response_id
())
;
evnt_vec
events
;
evnt_vec
events
;
args_vec
input_buffers
;
args_vec
input_buffers
;
args_vec
output_buffers
;
args_vec
output_buffers
;
size_vec
result_sizes
;
size_vec
result_sizes
;
add_kernel_arguments
(
events
,
input_buffers
,
output_buffers
,
add_kernel_arguments
(
events
,
input_buffers
,
output_buffers
,
result_sizes
,
content
,
indices
);
result_sizes
,
content
,
indices
);
auto
cmd
=
make_counted
<
command_type
>
(
hdl
,
this
,
auto
cmd
=
make_counted
<
command_type
>
(
std
::
move
(
hdl
),
actor_cast
<
strong_actor_ptr
>
(
this
),
std
::
move
(
events
),
std
::
move
(
events
),
std
::
move
(
input_buffers
),
std
::
move
(
input_buffers
),
std
::
move
(
output_buffers
),
std
::
move
(
output_buffers
),
...
@@ -151,21 +168,29 @@ public:
...
@@ -151,21 +168,29 @@ public:
cmd
->
enqueue
();
cmd
->
enqueue
();
}
}
actor_facade
(
const
program
&
prog
,
kernel_ptr
kernel
,
void
enqueue
(
mailbox_element_ptr
ptr
,
execution_unit
*
eu
)
override
{
const
spawn_config
&
config
,
CAF_ASSERT
(
ptr
!=
nullptr
);
CAF_LOG_TRACE
(
CAF_ARG
(
*
ptr
));
enqueue
(
ptr
->
sender
,
ptr
->
mid
,
ptr
->
msg
,
eu
);
}
actor_facade
(
actor_config
actor_cfg
,
const
program
&
prog
,
kernel_ptr
kernel
,
const
spawn_config
&
spawn_cfg
,
input_mapping
map_args
,
output_mapping
map_result
,
input_mapping
map_args
,
output_mapping
map_result
,
std
::
tuple
<
Ts
...
>
xs
)
std
::
tuple
<
Ts
...
>
xs
)
:
kernel_
(
kernel
),
:
monitorable_actor
(
actor_cfg
),
kernel_
(
kernel
),
program_
(
prog
.
program_
),
program_
(
prog
.
program_
),
context_
(
prog
.
context_
),
context_
(
prog
.
context_
),
queue_
(
prog
.
queue_
),
queue_
(
prog
.
queue_
),
config_
(
confi
g
),
spawn_cfg_
(
spawn_cf
g
),
map_args_
(
std
::
move
(
map_args
)),
map_args_
(
std
::
move
(
map_args
)),
map_results_
(
std
::
move
(
map_result
)),
map_results_
(
std
::
move
(
map_result
)),
argument_types_
(
xs
)
{
argument_types_
(
xs
)
{
CAF_LOG_TRACE
(
"id: "
<<
this
->
id
(
));
CAF_LOG_TRACE
(
CAF_ARG
(
this
->
id
()
));
default_output_size_
=
std
::
accumulate
(
confi
g_
.
dimensions
().
begin
(),
default_output_size_
=
std
::
accumulate
(
spawn_cf
g_
.
dimensions
().
begin
(),
confi
g_
.
dimensions
().
end
(),
spawn_cf
g_
.
dimensions
().
end
(),
size_t
{
1
},
size_t
{
1
},
std
::
multiplies
<
size_t
>
{});
std
::
multiplies
<
size_t
>
{});
}
}
...
@@ -258,7 +283,7 @@ public:
...
@@ -258,7 +283,7 @@ public:
program_ptr
program_
;
program_ptr
program_
;
context_ptr
context_
;
context_ptr
context_
;
command_queue_ptr
queue_
;
command_queue_ptr
queue_
;
spawn_config
confi
g_
;
spawn_config
spawn_cf
g_
;
input_mapping
map_args_
;
input_mapping
map_args_
;
output_mapping
map_results_
;
output_mapping
map_results_
;
std
::
tuple
<
Ts
...
>
argument_types_
;
std
::
tuple
<
Ts
...
>
argument_types_
;
...
...
libcaf_opencl/caf/opencl/all.hpp
View file @
108e51ff
...
@@ -21,6 +21,6 @@
...
@@ -21,6 +21,6 @@
#ifndef CAF_OPENCL_ALL_HPP
#ifndef CAF_OPENCL_ALL_HPP
#define CAF_OPENCL_ALL_HPP
#define CAF_OPENCL_ALL_HPP
#include "caf/opencl/
spawn_cl
.hpp"
#include "caf/opencl/
manager
.hpp"
#endif // CAF_OPENCL_ALL_HPP
#endif // CAF_OPENCL_ALL_HPP
libcaf_opencl/caf/opencl/arguments.hpp
View file @
108e51ff
...
@@ -25,7 +25,7 @@
...
@@ -25,7 +25,7 @@
#include <type_traits>
#include <type_traits>
#include "caf/message.hpp"
#include "caf/message.hpp"
#include "caf/
maybe
.hpp"
#include "caf/
optional
.hpp"
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
...
@@ -56,7 +56,7 @@ struct out {
...
@@ -56,7 +56,7 @@ struct out {
out
()
{
}
out
()
{
}
template
<
class
F
>
template
<
class
F
>
out
(
F
fun
)
{
out
(
F
fun
)
{
fun_
=
[
fun
](
message
&
msg
)
->
maybe
<
size_t
>
{
fun_
=
[
fun
](
message
&
msg
)
->
optional
<
size_t
>
{
auto
res
=
msg
.
apply
(
fun
);
auto
res
=
msg
.
apply
(
fun
);
size_t
result
;
size_t
result
;
if
(
res
)
{
if
(
res
)
{
...
@@ -66,10 +66,10 @@ struct out {
...
@@ -66,10 +66,10 @@ struct out {
return
none
;
return
none
;
};
};
}
}
maybe
<
size_t
>
operator
()(
message
&
msg
)
const
{
optional
<
size_t
>
operator
()(
message
&
msg
)
const
{
return
fun_
?
fun_
(
msg
)
:
0
;
return
fun_
?
fun_
(
msg
)
:
0
UL
;
}
}
std
::
function
<
maybe
<
size_t
>
(
message
&
)
>
fun_
;
std
::
function
<
optional
<
size_t
>
(
message
&
)
>
fun_
;
};
};
...
...
libcaf_opencl/caf/opencl/command.hpp
View file @
108e51ff
...
@@ -26,10 +26,11 @@
...
@@ -26,10 +26,11 @@
#include <algorithm>
#include <algorithm>
#include <functional>
#include <functional>
#include "caf/logger.hpp"
#include "caf/actor_cast.hpp"
#include "caf/abstract_actor.hpp"
#include "caf/abstract_actor.hpp"
#include "caf/response_promise.hpp"
#include "caf/response_promise.hpp"
#include "caf/detail/logging.hpp"
#include "caf/detail/scope_guard.hpp"
#include "caf/detail/scope_guard.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/global.hpp"
...
@@ -43,14 +44,14 @@ namespace opencl {
...
@@ -43,14 +44,14 @@ namespace opencl {
template
<
class
FacadeType
,
class
...
Ts
>
template
<
class
FacadeType
,
class
...
Ts
>
class
command
:
public
ref_counted
{
class
command
:
public
ref_counted
{
public:
public:
command
(
response_promise
handle
,
intrusive_ptr
<
FacadeType
>
actor_facade
,
command
(
std
::
tuple
<
strong_actor_ptr
,
message_id
>
handle
,
strong_actor_ptr
actor_facade
,
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
mem_ptr
>
input_buffers
,
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
mem_ptr
>
input_buffers
,
std
::
vector
<
mem_ptr
>
output_buffers
,
std
::
vector
<
size_t
>
result_sizes
,
std
::
vector
<
mem_ptr
>
output_buffers
,
std
::
vector
<
size_t
>
result_sizes
,
message
msg
)
message
msg
)
:
result_sizes_
(
result_sizes
),
:
result_sizes_
(
result_sizes
),
handle_
(
handle
),
handle_
(
handle
),
actor_facade_
(
actor_facade
),
actor_facade_
(
actor_facade
),
queue_
(
actor_facade
->
queue_
),
mem_in_events_
(
std
::
move
(
events
)),
mem_in_events_
(
std
::
move
(
events
)),
input_buffers_
(
std
::
move
(
input_buffers
)),
input_buffers_
(
std
::
move
(
input_buffers
)),
output_buffers_
(
std
::
move
(
output_buffers
)),
output_buffers_
(
std
::
move
(
output_buffers
)),
...
@@ -76,18 +77,21 @@ public:
...
@@ -76,18 +77,21 @@ public:
auto
data_or_nullptr
=
[](
const
dim_vec
&
vec
)
{
auto
data_or_nullptr
=
[](
const
dim_vec
&
vec
)
{
return
vec
.
empty
()
?
nullptr
:
vec
.
data
();
return
vec
.
empty
()
?
nullptr
:
vec
.
data
();
};
};
auto
actor_facade
=
static_cast
<
FacadeType
*>
(
actor_cast
<
abstract_actor
*>
(
actor_facade_
));
// OpenCL expects cl_uint (unsigned int), hence the cast
// OpenCL expects cl_uint (unsigned int), hence the cast
cl_int
err
=
clEnqueueNDRangeKernel
(
cl_int
err
=
clEnqueueNDRangeKernel
(
queue_
.
get
(),
actor_facade_
->
kernel_
.
get
(),
actor_facade
->
queue_
.
get
(),
actor_facade
->
kernel_
.
get
(),
static_cast
<
cl_uint
>
(
actor_facade
_
->
confi
g_
.
dimensions
().
size
()),
static_cast
<
cl_uint
>
(
actor_facade
->
spawn_cf
g_
.
dimensions
().
size
()),
data_or_nullptr
(
actor_facade
_
->
confi
g_
.
offsets
()),
data_or_nullptr
(
actor_facade
->
spawn_cf
g_
.
offsets
()),
data_or_nullptr
(
actor_facade
_
->
confi
g_
.
dimensions
()),
data_or_nullptr
(
actor_facade
->
spawn_cf
g_
.
dimensions
()),
data_or_nullptr
(
actor_facade
_
->
confi
g_
.
local_dimensions
()),
data_or_nullptr
(
actor_facade
->
spawn_cf
g_
.
local_dimensions
()),
static_cast
<
cl_uint
>
(
mem_in_events_
.
size
()),
static_cast
<
cl_uint
>
(
mem_in_events_
.
size
()),
(
mem_in_events_
.
empty
()
?
nullptr
:
mem_in_events_
.
data
()),
&
event_k
(
mem_in_events_
.
empty
()
?
nullptr
:
mem_in_events_
.
data
()),
&
event_k
);
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOGMF
(
CAF_ERROR
,
"clEnqueueNDRangeKernel: "
<<
get_opencl_error
(
err
));
CAF_LOG_ERROR
(
"clEnqueueNDRangeKernel: "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
clReleaseEvent
(
event_k
);
clReleaseEvent
(
event_k
);
this
->
deref
();
this
->
deref
();
return
;
return
;
...
@@ -96,15 +100,15 @@ public:
...
@@ -96,15 +100,15 @@ public:
cl_event
marker
;
cl_event
marker
;
#if defined(__APPLE__)
#if defined(__APPLE__)
err
=
clEnqueueMarkerWithWaitList
(
err
=
clEnqueueMarkerWithWaitList
(
queue_
.
get
(),
actor_facade
->
queue_
.
get
(),
static_cast
<
cl_uint
>
(
mem_out_events_
.
size
()),
static_cast
<
cl_uint
>
(
mem_out_events_
.
size
()),
mem_out_events_
.
data
(),
&
marker
mem_out_events_
.
data
(),
&
marker
);
);
#else
#else
err
=
clEnqueueMarker
(
queue_
.
get
(),
&
marker
);
err
=
clEnqueueMarker
(
actor_facade
->
queue_
.
get
(),
&
marker
);
#endif
#endif
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG
MF
(
CAF_ERROR
,
"clSetEventCallback: "
<<
get_opencl_error
(
err
));
CAF_LOG
_ERROR
(
"clSetEventCallback: "
<<
CAF_ARG
(
get_opencl_error
(
err
)
));
clReleaseEvent
(
marker
);
clReleaseEvent
(
marker
);
clReleaseEvent
(
event_k
);
clReleaseEvent
(
event_k
);
this
->
deref
();
// callback is not set
this
->
deref
();
// callback is not set
...
@@ -118,15 +122,15 @@ public:
...
@@ -118,15 +122,15 @@ public:
},
},
this
);
this
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG
MF
(
CAF_ERROR
,
"clSetEventCallback: "
<<
get_opencl_error
(
err
));
CAF_LOG
_ERROR
(
"clSetEventCallback: "
<<
CAF_ARG
(
get_opencl_error
(
err
)
));
clReleaseEvent
(
marker
);
clReleaseEvent
(
marker
);
clReleaseEvent
(
event_k
);
clReleaseEvent
(
event_k
);
this
->
deref
();
// callback is not set
this
->
deref
();
// callback is not set
return
;
return
;
}
}
err
=
clFlush
(
queue_
.
get
());
err
=
clFlush
(
actor_facade
->
queue_
.
get
());
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG
MF
(
CAF_ERROR
,
"clFlush: "
<<
get_opencl_error
(
err
));
CAF_LOG
_ERROR
(
"clFlush: "
<<
CAF_ARG
(
get_opencl_error
(
err
)
));
}
}
mem_out_events_
.
push_back
(
std
::
move
(
event_k
));
mem_out_events_
.
push_back
(
std
::
move
(
event_k
));
mem_out_events_
.
push_back
(
std
::
move
(
marker
));
mem_out_events_
.
push_back
(
std
::
move
(
marker
));
...
@@ -135,9 +139,8 @@ public:
...
@@ -135,9 +139,8 @@ public:
private:
private:
std
::
vector
<
size_t
>
result_sizes_
;
std
::
vector
<
size_t
>
result_sizes_
;
response_promise
handle_
;
std
::
tuple
<
strong_actor_ptr
,
message_id
>
handle_
;
intrusive_ptr
<
FacadeType
>
actor_facade_
;
strong_actor_ptr
actor_facade_
;
command_queue_ptr
queue_
;
std
::
vector
<
cl_event
>
mem_in_events_
;
std
::
vector
<
cl_event
>
mem_in_events_
;
std
::
vector
<
cl_event
>
mem_out_events_
;
std
::
vector
<
cl_event
>
mem_out_events_
;
std
::
vector
<
mem_ptr
>
input_buffers_
;
std
::
vector
<
mem_ptr
>
input_buffers_
;
...
@@ -151,6 +154,8 @@ private:
...
@@ -151,6 +154,8 @@ private:
template
<
long
I
,
long
...
Is
>
template
<
long
I
,
long
...
Is
>
void
enqueue_read_buffers
(
cl_event
&
kernel_done
,
detail
::
int_list
<
I
,
Is
...
>
)
{
void
enqueue_read_buffers
(
cl_event
&
kernel_done
,
detail
::
int_list
<
I
,
Is
...
>
)
{
auto
actor_facade
=
static_cast
<
FacadeType
*>
(
actor_cast
<
abstract_actor
*>
(
actor_facade_
));
using
container_type
=
using
container_type
=
typename
std
::
tuple_element
<
I
,
std
::
tuple
<
Ts
...
>>::
type
;
typename
std
::
tuple_element
<
I
,
std
::
tuple
<
Ts
...
>>::
type
;
using
value_type
=
typename
container_type
::
value_type
;
using
value_type
=
typename
container_type
::
value_type
;
...
@@ -158,7 +163,8 @@ private:
...
@@ -158,7 +163,8 @@ private:
auto
size
=
result_sizes_
[
I
];
auto
size
=
result_sizes_
[
I
];
auto
buffer_size
=
sizeof
(
value_type
)
*
result_sizes_
[
I
];
auto
buffer_size
=
sizeof
(
value_type
)
*
result_sizes_
[
I
];
std
::
get
<
I
>
(
result_buffers_
).
resize
(
size
);
std
::
get
<
I
>
(
result_buffers_
).
resize
(
size
);
auto
err
=
clEnqueueReadBuffer
(
queue_
.
get
(),
output_buffers_
[
I
].
get
(),
auto
err
=
clEnqueueReadBuffer
(
actor_facade
->
queue_
.
get
(),
output_buffers_
[
I
].
get
(),
CL_FALSE
,
0
,
buffer_size
,
CL_FALSE
,
0
,
buffer_size
,
std
::
get
<
I
>
(
result_buffers_
).
data
(),
std
::
get
<
I
>
(
result_buffers_
).
data
(),
1
,
&
kernel_done
,
&
event
);
1
,
&
kernel_done
,
&
event
);
...
@@ -172,12 +178,15 @@ private:
...
@@ -172,12 +178,15 @@ private:
}
}
void
handle_results
()
{
void
handle_results
()
{
auto
&
map_fun
=
actor_facade_
->
map_results_
;
auto
actor_facade
=
static_cast
<
FacadeType
*>
(
actor_cast
<
abstract_actor
*>
(
actor_facade_
));
auto
&
map_fun
=
actor_facade
->
map_results_
;
auto
msg
=
map_fun
?
apply_args
(
map_fun
,
auto
msg
=
map_fun
?
apply_args
(
map_fun
,
detail
::
get_indices
(
result_buffers_
),
detail
::
get_indices
(
result_buffers_
),
result_buffers_
)
result_buffers_
)
:
message_from_results
{}(
result_buffers_
);
:
message_from_results
{}(
result_buffers_
);
handle_
.
deliver
(
std
::
move
(
msg
));
get
<
0
>
(
handle_
)
->
enqueue
(
actor_facade_
,
get
<
1
>
(
handle_
),
std
::
move
(
msg
),
nullptr
);
}
}
};
};
...
...
libcaf_opencl/caf/opencl/
metainfo
.hpp
→
libcaf_opencl/caf/opencl/
detail/spawn_helper
.hpp
View file @
108e51ff
...
@@ -17,68 +17,45 @@
...
@@ -17,68 +17,45 @@
* http://www.boost.org/LICENSE_1_0.txt. *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
******************************************************************************/
#ifndef CAF_
METAINFO
_HPP
#ifndef CAF_
OPENCL_DETAIL_SPAWN_HELPER
_HPP
#define CAF_
METAINFO
_HPP
#define CAF_
OPENCL_DETAIL_SPAWN_HELPER
_HPP
#include <atomic>
#include <vector>
#include <algorithm>
#include <functional>
#include "caf/all.hpp"
#include "caf/config.hpp"
#include "caf/maybe.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/actor_facade.hpp"
#include "caf/opencl/actor_facade.hpp"
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
namespace
detail
{
class
metainfo
:
public
detail
::
abstract_singleton
{
struct
tuple_construct
{
};
friend
class
program
;
friend
class
detail
::
singletons
;
template
<
class
...
Ts
>
friend
command_queue_ptr
get_command_queue
(
uint32_t
id
);
struct
cl_spawn_helper
{
using
impl
=
opencl
::
actor_facade
<
Ts
...
>
;
public:
using
map_in_fun
=
std
::
function
<
optional
<
message
>
(
message
&
)
>
;
/// Get a list of all available devices. This is depricated, use the more specific
using
map_out_fun
=
typename
impl
::
output_mapping
;
/// get_deivce and get_deivce_if functions.
/// (Returns only devices of the first discovered platform).
actor
operator
()(
actor_config
actor_cfg
,
const
opencl
::
program
&
p
,
const
std
::
vector
<
device
>&
get_devices
()
const
CAF_DEPRECATED
;
const
char
*
fn
,
const
opencl
::
spawn_config
&
spawn_cfg
,
/// Get the device with id. These ids are assigned sequientally to all available devices.
Ts
&&
...
xs
)
const
{
const
maybe
<
const
device
&>
get_device
(
size_t
id
=
0
)
const
;
return
actor_cast
<
actor
>
(
impl
::
create
(
std
::
move
(
actor_cfg
),
/// Get the first device that satisfies the predicate.
p
,
fn
,
spawn_cfg
,
/// The predicate should accept a `const device&` and return a bool;
map_in_fun
{},
map_out_fun
{},
template
<
class
UnaryPredicate
>
std
::
move
(
xs
)...));
const
maybe
<
const
device
&>
get_device_if
(
UnaryPredicate
p
)
const
{
for
(
auto
&
pl
:
platforms_
)
{
for
(
auto
&
dev
:
pl
.
get_devices
())
{
if
(
p
(
dev
))
return
dev
;
}
}
actor
operator
()(
actor_config
actor_cfg
,
const
opencl
::
program
&
p
,
const
char
*
fn
,
const
opencl
::
spawn_config
&
spawn_cfg
,
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
,
spawn_cfg
,
std
::
move
(
map_input
),
std
::
move
(
map_output
),
std
::
move
(
xs
)...));
}
}
return
none
;
}
/// Get metainfo instance.
static
metainfo
*
instance
();
private:
metainfo
()
=
default
;
void
stop
()
override
;
void
initialize
()
override
;
void
dispose
()
override
;
std
::
vector
<
platform
>
platforms_
;
};
};
}
// namespace detail
}
// namespace opencl
}
// namespace opencl
}
// namespace caf
}
// namespace caf
#endif // CAF_METAINFO
_HPP
#endif // CAF_OPENCL_DETAIL_SPAWN_HELPER
_HPP
libcaf_opencl/caf/opencl/device.hpp
View file @
108e51ff
...
@@ -29,11 +29,13 @@ namespace caf {
...
@@ -29,11 +29,13 @@ namespace caf {
namespace
opencl
{
namespace
opencl
{
class
program
;
class
program
;
class
manager
;
class
device
{
class
device
{
public:
friend
class
program
;
friend
class
program
;
friend
class
manager
;
public:
/// Intialize a new device in a context using a sepcific device_id
/// Intialize a new device in a context using a sepcific device_id
static
device
create
(
context_ptr
context
,
device_ptr
device_id
,
unsigned
id
);
static
device
create
(
context_ptr
context
,
device_ptr
device_id
,
unsigned
id
);
/// Get the id assigned by caf
/// Get the id assigned by caf
...
...
libcaf_opencl/caf/opencl/manager.hpp
0 → 100644
View file @
108e51ff
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2015 *
* Dominik Charousset <dominik.charousset (at) haw-hamburg.de> *
* *
* Distributed under the terms and conditions of the BSD 3-Clause License or *
* (at your option) under the terms and conditions of the Boost Software *
* License 1.0. See accompanying files LICENSE and 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. *
******************************************************************************/
#ifndef CAF_OPENCL_MANAGER_HPP
#define CAF_OPENCL_MANAGER_HPP
#include <atomic>
#include <vector>
#include <algorithm>
#include <functional>
#include "caf/optional.hpp"
#include "caf/config.hpp"
#include "caf/actor_system.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/actor_facade.hpp"
#include "caf/opencl/detail/spawn_helper.hpp"
namespace
caf
{
namespace
opencl
{
class
manager
:
public
actor_system
::
module
{
public:
friend
class
program
;
friend
class
actor_system
;
friend
command_queue_ptr
get_command_queue
(
uint32_t
id
);
manager
(
const
manager
&
)
=
delete
;
manager
&
operator
=
(
const
manager
&
)
=
delete
;
/// Get the device with id, which is assigned sequientally.
const
optional
<
const
device
&>
get_device
(
size_t
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
>
const
optional
<
const
device
&>
get_device_if
(
UnaryPredicate
p
)
const
{
for
(
auto
&
pl
:
platforms_
)
{
for
(
auto
&
dev
:
pl
.
get_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
,
caf
::
detail
::
type_list
<>
);
// OpenCL functionality
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
program
create_program
(
const
char
*
kernel_source
,
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
create_program
(
const
char
*
kernel_source
,
const
char
*
options
,
const
device
&
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
>
typename
std
::
enable_if
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>::
type
spawn
(
const
opencl
::
program
&
prog
,
const
char
*
fname
,
const
opencl
::
spawn_config
&
config
,
T
x
,
Ts
...
xs
)
{
detail
::
cl_spawn_helper
<
T
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
prog
,
fname
,
config
,
std
::
move
(
x
),
std
::
move
(
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
>
typename
std
::
enable_if
<
opencl
::
is_opencl_arg
<
T
>::
value
,
actor
>::
type
spawn
(
const
char
*
source
,
const
char
*
fname
,
const
opencl
::
spawn_config
&
config
,
T
x
,
Ts
...
xs
)
{
detail
::
cl_spawn_helper
<
T
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
config
,
std
::
move
(
x
),
std
::
move
(
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
&
prog
,
const
char
*
fname
,
const
opencl
::
spawn_config
&
config
,
std
::
function
<
optional
<
message
>
(
message
&
)
>
map_args
,
Fun
map_result
,
Ts
...
xs
)
{
detail
::
cl_spawn_helper
<
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
prog
,
fname
,
config
,
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
::
spawn_config
&
config
,
std
::
function
<
optional
<
message
>
(
message
&
)
>
map_args
,
Fun
map_result
,
Ts
...
xs
)
{
detail
::
cl_spawn_helper
<
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
config
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
std
::
forward
<
Ts
>
(
xs
)...);
}
protected:
manager
(
actor_system
&
sys
);
~
manager
();
private:
actor_system
&
system_
;
std
::
vector
<
platform
>
platforms_
;
};
}
// namespace opencl
}
// namespace caf
#endif // CAF_OPENCL_MANAGER_HPP
libcaf_opencl/caf/opencl/opencl_err.hpp
View file @
108e51ff
...
@@ -24,7 +24,7 @@
...
@@ -24,7 +24,7 @@
#include "caf/opencl/global.hpp"
#include "caf/opencl/global.hpp"
#include "caf/
detail/logging
.hpp"
#include "caf/
logger
.hpp"
#define CAF_CLF(funname) #funname , funname
#define CAF_CLF(funname) #funname , funname
...
...
libcaf_opencl/caf/opencl/platform.hpp
View file @
108e51ff
...
@@ -26,10 +26,9 @@ namespace caf {
...
@@ -26,10 +26,9 @@ namespace caf {
namespace
opencl
{
namespace
opencl
{
class
platform
{
class
platform
{
public:
friend
class
program
;
friend
class
program
;
public:
inline
const
std
::
vector
<
device
>&
get_devices
()
const
;
inline
const
std
::
vector
<
device
>&
get_devices
()
const
;
inline
const
std
::
string
&
get_name
()
const
;
inline
const
std
::
string
&
get_name
()
const
;
inline
const
std
::
string
&
get_vendor
()
const
;
inline
const
std
::
string
&
get_vendor
()
const
;
...
...
libcaf_opencl/caf/opencl/program.hpp
View file @
108e51ff
...
@@ -35,29 +35,19 @@ class actor_facade;
...
@@ -35,29 +35,19 @@ class actor_facade;
/// @brief A wrapper for OpenCL's cl_program.
/// @brief A wrapper for OpenCL's cl_program.
class
program
{
class
program
{
public:
friend
class
manager
;
template
<
class
...
Ts
>
template
<
class
...
Ts
>
friend
class
actor_facade
;
friend
class
actor_facade
;
public:
/// @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
/// @returns A program object.
static
program
create
(
const
char
*
kernel_source
,
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.
static
program
create
(
const
char
*
kernel_source
,
const
char
*
options
,
const
device
&
dev
);
private:
private:
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program_ptr
prog
);
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program_ptr
prog
,
std
::
map
<
std
::
string
,
kernel_ptr
>
available_kernels
);
context_ptr
context_
;
context_ptr
context_
;
program_ptr
program_
;
program_ptr
program_
;
command_queue_ptr
queue_
;
command_queue_ptr
queue_
;
// save CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE
std
::
map
<
std
::
string
,
kernel_ptr
>
available_kernels_
;
};
};
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/caf/opencl/spawn_cl.hpp
deleted
100644 → 0
View file @
785153d3
This diff is collapsed.
Click to expand it.
libcaf_opencl/examples/proper_matrix.cpp
View file @
108e51ff
...
@@ -24,13 +24,15 @@
...
@@ -24,13 +24,15 @@
#include <iostream>
#include <iostream>
#include "caf/all.hpp"
#include "caf/all.hpp"
#include "caf/opencl/spawn_cl.hpp"
#include "caf/detail/limited_vector.hpp"
#include "caf/opencl/all.hpp"
using
namespace
std
;
using
namespace
std
;
using
namespace
caf
;
using
namespace
caf
;
using
namespace
caf
::
opencl
;
using
namespace
caf
::
opencl
;
using
detail
::
limited_vector
;
using
caf
::
detail
::
limited_vector
;
namespace
{
namespace
{
...
@@ -61,11 +63,16 @@ constexpr const char* kernel_source = R"__(
...
@@ -61,11 +63,16 @@ constexpr const char* kernel_source = R"__(
template
<
size_t
Size
>
template
<
size_t
Size
>
class
square_matrix
{
class
square_matrix
{
public:
public:
using
value_type
=
fvec
::
value_type
;
static
constexpr
size_t
num_elements
=
Size
*
Size
;
static
constexpr
size_t
num_elements
=
Size
*
Size
;
// allows serialization
template
<
class
IO
>
friend
void
serialize
(
IO
&
in_or_out
,
square_matrix
&
m
,
const
unsigned
int
)
{
in_or_out
&
m
.
data_
;
}
square_matrix
(
square_matrix
&&
)
=
default
;
square_matrix
(
square_matrix
&&
)
=
default
;
square_matrix
(
const
square_matrix
&
)
=
default
;
square_matrix
(
const
square_matrix
&
)
=
default
;
square_matrix
&
operator
=
(
square_matrix
&&
)
=
default
;
square_matrix
&
operator
=
(
square_matrix
&&
)
=
default
;
...
@@ -101,9 +108,7 @@ public:
...
@@ -101,9 +108,7 @@ public:
const
fvec
&
data
()
const
{
return
data_
;
}
const
fvec
&
data
()
const
{
return
data_
;
}
private:
private:
fvec
data_
;
fvec
data_
;
};
};
template
<
size_t
Size
>
template
<
size_t
Size
>
...
@@ -135,6 +140,7 @@ inline bool operator!=(const square_matrix<Size>& lhs,
...
@@ -135,6 +140,7 @@ inline bool operator!=(const square_matrix<Size>& lhs,
using
matrix_type
=
square_matrix
<
matrix_size
>
;
using
matrix_type
=
square_matrix
<
matrix_size
>
;
void
multiplier
(
event_based_actor
*
self
)
{
void
multiplier
(
event_based_actor
*
self
)
{
auto
&
mngr
=
self
->
system
().
opencl_manager
();
// create two matrices with ascending values
// create two matrices with ascending values
matrix_type
m1
;
matrix_type
m1
;
...
@@ -145,7 +151,7 @@ void multiplier(event_based_actor* self) {
...
@@ -145,7 +151,7 @@ void multiplier(event_based_actor* self) {
cout
<<
"calculating square of matrix:"
<<
endl
cout
<<
"calculating square of matrix:"
<<
endl
<<
to_string
(
m1
)
<<
endl
;
<<
to_string
(
m1
)
<<
endl
;
auto
unbox_args
=
[](
message
&
msg
)
->
maybe
<
message
>
{
auto
unbox_args
=
[](
message
&
msg
)
->
optional
<
message
>
{
return
msg
.
apply
(
return
msg
.
apply
(
[](
matrix_type
&
lhs
,
matrix_type
&
rhs
)
{
[](
matrix_type
&
lhs
,
matrix_type
&
rhs
)
{
return
make_message
(
std
::
move
(
lhs
.
data
()),
std
::
move
(
rhs
.
data
()));
return
make_message
(
std
::
move
(
lhs
.
data
()),
std
::
move
(
rhs
.
data
()));
...
@@ -171,14 +177,14 @@ void multiplier(event_based_actor* self) {
...
@@ -171,14 +177,14 @@ void multiplier(event_based_actor* self) {
// used as response message
// used as response message
// from 6 : a description of the kernel signature using in/out/in_out classes
// from 6 : a description of the kernel signature using in/out/in_out classes
// with the argument type packed in vectors
// with the argument type packed in vectors
auto
worker
=
spawn_cl
(
kernel_source
,
kernel_name
,
auto
worker
=
mngr
.
spawn
(
kernel_source
,
kernel_name
,
spawn_config
{
dim_vec
{
matrix_size
,
matrix_size
}},
spawn_config
{
dim_vec
{
matrix_size
,
matrix_size
}},
unbox_args
,
box_res
,
unbox_args
,
box_res
,
in
<
fvec
>
{},
in
<
fvec
>
{},
out
<
fvec
>
{});
in
<
fvec
>
{},
in
<
fvec
>
{},
out
<
fvec
>
{});
// send both matrices to the actor and
// send both matrices to the actor and
// wait for results in form of a matrix_type
// wait for results in form of a matrix_type
self
->
sync_send
(
worker
,
move
(
m1
),
move
(
m2
)).
then
(
self
->
request
(
worker
,
chrono
::
seconds
(
5
)
,
move
(
m1
),
move
(
m2
)).
then
(
[](
const
matrix_type
&
result
)
{
[](
const
matrix_type
&
result
)
{
cout
<<
"result:"
<<
endl
<<
to_string
(
result
);
cout
<<
"result:"
<<
endl
<<
to_string
(
result
);
}
}
...
@@ -188,9 +194,12 @@ void multiplier(event_based_actor* self) {
...
@@ -188,9 +194,12 @@ void multiplier(event_based_actor* self) {
int
main
()
{
int
main
()
{
// matrix_type ist not a simple type,
// matrix_type ist not a simple type,
// it must be annouced to libcaf
// it must be annouced to libcaf
announce
<
matrix_type
>
(
"matrix_type"
);
actor_system_config
cfg
;
spawn
(
multiplier
);
cfg
.
load
<
opencl
::
manager
>
()
await_all_actors_done
();
.
add_message_type
<
fvec
>
(
"float_vector"
)
shutdown
();
.
add_message_type
<
matrix_type
>
(
"square_matrix"
);
actor_system
system
{
cfg
};
system
.
spawn
(
multiplier
);
system
.
await_all_actors_done
();
return
0
;
return
0
;
}
}
libcaf_opencl/examples/simple_matrix.cpp
View file @
108e51ff
...
@@ -23,13 +23,13 @@
...
@@ -23,13 +23,13 @@
#include <iostream>
#include <iostream>
#include "caf/all.hpp"
#include "caf/all.hpp"
#include "caf/opencl/
spawn_c
l.hpp"
#include "caf/opencl/
al
l.hpp"
using
namespace
std
;
using
namespace
std
;
using
namespace
caf
;
using
namespace
caf
;
using
namespace
caf
::
opencl
;
using
namespace
caf
::
opencl
;
using
detail
::
limited_vector
;
using
caf
::
detail
::
limited_vector
;
namespace
{
namespace
{
...
@@ -93,11 +93,13 @@ void multiplier(event_based_actor* self) {
...
@@ -93,11 +93,13 @@ void multiplier(event_based_actor* self) {
// - local dimensions (optional)
// - local dimensions (optional)
// 4th to Nth arg: the kernel signature described by in/out/in_out classes
// 4th to Nth arg: the kernel signature described by in/out/in_out classes
// that contain the argument type in their template, requires vectors
// that contain the argument type in their template, requires vectors
auto
worker
=
spawn_cl
(
kernel_source
,
kernel_name
,
auto
worker
=
self
->
system
().
opencl_manager
().
spawn
(
kernel_source
,
kernel_name
,
spawn_config
{
dim_vec
{
matrix_size
,
matrix_size
}},
spawn_config
{
dim_vec
{
matrix_size
,
matrix_size
}},
in
<
fvec
>
{},
in
<
fvec
>
{},
out
<
fvec
>
{});
in
<
fvec
>
{},
in
<
fvec
>
{},
out
<
fvec
>
{}
);
// send both matrices to the actor and wait for a result
// send both matrices to the actor and wait for a result
self
->
sync_send
(
worker
,
move
(
m1
),
move
(
m2
)).
then
(
self
->
request
(
worker
,
chrono
::
seconds
(
5
)
,
move
(
m1
),
move
(
m2
)).
then
(
[](
const
fvec
&
result
)
{
[](
const
fvec
&
result
)
{
cout
<<
"result: "
<<
endl
;
cout
<<
"result: "
<<
endl
;
print_as_matrix
(
result
);
print_as_matrix
(
result
);
...
@@ -106,9 +108,11 @@ void multiplier(event_based_actor* self) {
...
@@ -106,9 +108,11 @@ void multiplier(event_based_actor* self) {
}
}
int
main
()
{
int
main
()
{
announce
<
fvec
>
(
"float_vector"
);
actor_system_config
cfg
;
spawn
(
multiplier
);
cfg
.
load
<
opencl
::
manager
>
()
await_all_actors_done
();
.
add_message_type
<
fvec
>
(
"float_vector"
);
shutdown
();
actor_system
system
{
cfg
};
system
.
spawn
(
multiplier
);
system
.
await_all_actors_done
();
return
0
;
return
0
;
}
}
libcaf_opencl/src/device.cpp
View file @
108e51ff
...
@@ -19,6 +19,7 @@
...
@@ -19,6 +19,7 @@
#include <iostream>
#include <iostream>
#include "caf/logger.hpp"
#include "caf/string_algorithms.hpp"
#include "caf/string_algorithms.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/global.hpp"
...
@@ -31,8 +32,7 @@ namespace caf {
...
@@ -31,8 +32,7 @@ namespace caf {
namespace
opencl
{
namespace
opencl
{
device
device
::
create
(
context_ptr
context
,
device_ptr
device_id
,
unsigned
id
)
{
device
device
::
create
(
context_ptr
context
,
device_ptr
device_id
,
unsigned
id
)
{
CAF_LOGF_DEBUG
(
"creating device for opencl device id '"
CAF_LOG_DEBUG
(
"creating device for opencl device with id:"
<<
CAF_ARG
(
id
));
<<
device_id
.
get
()
<<
"' with id '"
<<
id
<<
"'"
);
// look up properties we need to create the command queue
// look up properties we need to create the command queue
auto
supported
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_QUEUE_PROPERTIES
);
auto
supported
=
info
<
cl_ulong
>
(
device_id
,
CL_DEVICE_QUEUE_PROPERTIES
);
bool
profiling
=
supported
&
CL_QUEUE_PROFILING_ENABLE
;
bool
profiling
=
supported
&
CL_QUEUE_PROFILING_ENABLE
;
...
...
libcaf_opencl/src/m
etainfo
.cpp
→
libcaf_opencl/src/m
anager
.cpp
View file @
108e51ff
...
@@ -18,26 +18,19 @@
...
@@ -18,26 +18,19 @@
* http://www.boost.org/LICENSE_1_0.txt. *
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
******************************************************************************/
#include "caf/detail/type_list.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/device.hpp"
#include "caf/opencl/m
etainfo
.hpp"
#include "caf/opencl/m
anager
.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/opencl_err.hpp"
using
namespace
std
;
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
metainfo
*
metainfo
::
instance
()
{
const
optional
<
const
device
&>
manager
::
get_device
(
size_t
id
)
const
{
auto
sid
=
detail
::
singletons
::
opencl_plugin_id
;
auto
fac
=
[]
{
return
new
metainfo
;
};
auto
res
=
detail
::
singletons
::
get_plugin_singleton
(
sid
,
fac
);
return
static_cast
<
metainfo
*>
(
res
);
}
const
std
::
vector
<
device
>&
metainfo
::
get_devices
()
const
{
return
platforms_
.
front
().
get_devices
();
}
const
maybe
<
const
device
&>
metainfo
::
get_device
(
size_t
id
)
const
{
if
(
platforms_
.
empty
())
if
(
platforms_
.
empty
())
return
none
;
return
none
;
size_t
to
=
0
;
size_t
to
=
0
;
...
@@ -50,7 +43,7 @@ const maybe<const device&> metainfo::get_device(size_t id) const{
...
@@ -50,7 +43,7 @@ const maybe<const device&> metainfo::get_device(size_t id) const{
return
none
;
return
none
;
}
}
void
m
etainfo
::
initialize
(
)
{
void
m
anager
::
init
(
actor_system_config
&
)
{
// get number of available platforms
// get number of available platforms
auto
num_platforms
=
v1get
<
cl_uint
>
(
CAF_CLF
(
clGetPlatformIDs
));
auto
num_platforms
=
v1get
<
cl_uint
>
(
CAF_CLF
(
clGetPlatformIDs
));
// get platform ids
// get platform ids
...
@@ -67,11 +60,114 @@ void metainfo::initialize() {
...
@@ -67,11 +60,114 @@ void metainfo::initialize() {
}
}
}
}
void
metainfo
::
dispose
()
{
void
manager
::
start
()
{
delete
this
;
// 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
manager
::
create_program
(
const
char
*
kernel_source
,
const
char
*
options
,
uint32_t
device_id
)
{
auto
dev
=
get_device
(
device_id
);
if
(
!
dev
)
{
ostringstream
oss
;
oss
<<
"No device with id '"
<<
device_id
<<
"' found."
;
CAF_LOG_ERROR
(
CAF_ARG
(
oss
.
str
()));
throw
runtime_error
(
oss
.
str
());
}
return
create_program
(
kernel_source
,
options
,
*
dev
);
}
program
manager
::
create_program
(
const
char
*
kernel_source
,
const
char
*
options
,
const
device
&
dev
)
{
// create program object from kernel source
size_t
kernel_source_length
=
strlen
(
kernel_source
);
program_ptr
pptr
;
auto
rawptr
=
v2get
(
CAF_CLF
(
clCreateProgramWithSource
),
dev
.
context_
.
get
(),
cl_uint
{
1
},
&
kernel_source
,
&
kernel_source_length
);
pptr
.
reset
(
rawptr
,
false
);
// build programm from program object
auto
dev_tmp
=
dev
.
device_id_
.
get
();
cl_int
err
=
clBuildProgram
(
pptr
.
get
(),
1
,
&
dev_tmp
,
options
,
nullptr
,
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
oss
<<
"clBuildProgram: "
<<
get_opencl_error
(
err
);
// the build log will be printed by the pfn_notify (see opencl/manger.cpp)
#ifndef __APPLE__
// seems that just apple implemented the
// pfn_notify callback, but we can get
// the build log
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
,
sizeof
(
buildlog_buffer_size
),
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
(
buffer
[
0
])
*
buildlog_buffer_size
,
buffer
.
data
(),
nullptr
);
ostringstream
ss
;
ss
<<
"Build log:
\n
"
<<
string
(
buffer
.
data
())
<<
"
\n
#######################################"
;
CAF_LOG_ERROR
(
CAF_ARG
(
ss
.
str
()));
}
#endif
throw
runtime_error
(
oss
.
str
());
}
map
<
string
,
kernel_ptr
>
available_kernels
;
cl_uint
number_of_kernels
=
0
;
clCreateKernelsInProgram
(
pptr
.
get
(),
0
,
nullptr
,
&
number_of_kernels
);
vector
<
cl_kernel
>
kernels
(
number_of_kernels
);
err
=
clCreateKernelsInProgram
(
pptr
.
get
(),
number_of_kernels
,
kernels
.
data
(),
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
oss
<<
"clCreateKernelsInProgram: "
<<
get_opencl_error
(
err
);
throw
runtime_error
(
oss
.
str
());
}
else
{
for
(
cl_uint
i
=
0
;
i
<
number_of_kernels
;
++
i
)
{
size_t
ret_size
;
clGetKernelInfo
(
kernels
[
i
],
CL_KERNEL_FUNCTION_NAME
,
0
,
nullptr
,
&
ret_size
);
vector
<
char
>
name
(
ret_size
);
err
=
clGetKernelInfo
(
kernels
[
i
],
CL_KERNEL_FUNCTION_NAME
,
ret_size
,
reinterpret_cast
<
void
*>
(
name
.
data
()),
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
oss
<<
"clGetKernelInfo (CL_KERNEL_FUNCTION_NAME): "
<<
get_opencl_error
(
err
);
throw
runtime_error
(
oss
.
str
());
}
kernel_ptr
kernel
;
kernel
.
reset
(
move
(
kernels
[
i
]));
available_kernels
.
emplace
(
string
(
name
.
data
()),
move
(
kernel
));
}
}
return
{
dev
.
context_
,
dev
.
command_queue_
,
pptr
,
move
(
available_kernels
)};
}
manager
::
manager
(
actor_system
&
sys
)
:
system_
(
sys
){
// nop
}
}
void
metainfo
::
stop
()
{
manager
::~
manager
()
{
// nop
// nop
}
}
...
...
libcaf_opencl/src/opencl_err.cpp
View file @
108e51ff
...
@@ -33,9 +33,10 @@ void throwcl(const char* fname, cl_int err) {
...
@@ -33,9 +33,10 @@ void throwcl(const char* fname, cl_int err) {
}
}
void
pfn_notify
(
const
char
*
errinfo
,
const
void
*
,
size_t
,
void
*
)
{
void
pfn_notify
(
const
char
*
errinfo
,
const
void
*
,
size_t
,
void
*
)
{
CAF_LOG
F
_ERROR
(
"
\n
##### Error message via pfn_notify #####
\n
"
CAF_LOG_ERROR
(
"
\n
##### Error message via pfn_notify #####
\n
"
<<
errinfo
<<
<<
errinfo
<<
"
\n
########################################"
);
"
\n
########################################"
);
static_cast
<
void
>
(
errinfo
);
// remove warning
}
}
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/src/platform.cpp
View file @
108e51ff
...
@@ -63,7 +63,7 @@ platform platform::create(cl_platform_id platform_id,
...
@@ -63,7 +63,7 @@ platform platform::create(cl_platform_id platform_id,
}
}
if
(
device_information
.
empty
())
{
if
(
device_information
.
empty
())
{
string
errstr
=
"no devices for the platform found"
;
string
errstr
=
"no devices for the platform found"
;
CAF_LOG
F_ERROR
(
errstr
);
CAF_LOG
_ERROR
(
CAF_ARG
(
errstr
)
);
throw
runtime_error
(
move
(
errstr
));
throw
runtime_error
(
move
(
errstr
));
}
}
auto
name
=
platform_info
(
platform_id
,
CL_PLATFORM_NAME
);
auto
name
=
platform_info
(
platform_id
,
CL_PLATFORM_NAME
);
...
...
libcaf_opencl/src/program.cpp
View file @
108e51ff
...
@@ -23,8 +23,8 @@
...
@@ -23,8 +23,8 @@
#include <cstring>
#include <cstring>
#include <iostream>
#include <iostream>
#include "caf/opencl/manager.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/metainfo.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/opencl_err.hpp"
using
namespace
std
;
using
namespace
std
;
...
@@ -32,63 +32,13 @@ using namespace std;
...
@@ -32,63 +32,13 @@ using namespace std;
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
program
::
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program_ptr
prog
)
program
::
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program_ptr
prog
,
map
<
string
,
kernel_ptr
>
available_kernels
)
:
context_
(
move
(
context
)),
:
context_
(
move
(
context
)),
program_
(
move
(
prog
)),
program_
(
move
(
prog
)),
queue_
(
move
(
queue
))
{}
queue_
(
move
(
queue
)),
available_kernels_
(
move
(
available_kernels
))
{
program
program
::
create
(
const
char
*
kernel_source
,
const
char
*
options
,
// nop
uint32_t
device_id
)
{
auto
info
=
metainfo
::
instance
();
auto
dev
=
info
->
get_device
(
device_id
);
if
(
!
dev
)
{
ostringstream
oss
;
oss
<<
"No device with id '"
<<
device_id
<<
"' found."
;
CAF_LOGF_ERROR
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
return
program
::
create
(
kernel_source
,
options
,
*
dev
);
}
program
program
::
create
(
const
char
*
kernel_source
,
const
char
*
options
,
const
device
&
dev
)
{
// create program object from kernel source
size_t
kernel_source_length
=
strlen
(
kernel_source
);
program_ptr
pptr
;
auto
rawptr
=
v2get
(
CAF_CLF
(
clCreateProgramWithSource
),
dev
.
context_
.
get
(),
cl_uint
{
1
},
&
kernel_source
,
&
kernel_source_length
);
pptr
.
reset
(
rawptr
,
false
);
// build programm from program object
auto
dev_tmp
=
dev
.
device_id_
.
get
();
cl_int
err
=
clBuildProgram
(
pptr
.
get
(),
1
,
&
dev_tmp
,
options
,
nullptr
,
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
oss
<<
"clBuildProgram: "
<<
get_opencl_error
(
err
);
// the build log will be printed by the
// pfn_notify (see metainfo.cpp)
#ifndef __APPLE__
// seems that just apple implemented the
// pfn_notify callback, but we can get
// the build log
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
,
sizeof
(
buildlog_buffer_size
),
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
(
buffer
[
0
])
*
buildlog_buffer_size
,
buffer
.
data
(),
nullptr
);
CAF_LOGF_ERROR
(
"Build log:
\n
"
+
string
(
buffer
.
data
())
+
"
\n
########################################"
);
}
#endif
throw
runtime_error
(
oss
.
str
());
}
return
{
dev
.
context_
,
dev
.
command_queue_
,
pptr
};
}
}
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/test/opencl.cpp
View file @
108e51ff
This diff is collapsed.
Click to expand it.
libcaf_opencl/test/opencl_deprecated.cpp
deleted
100644 → 0
View file @
785153d3
#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/opencl/spawn_cl.hpp"
CAF_PUSH_NO_DEPRECATED_WARNING
using
namespace
caf
;
using
namespace
caf
::
opencl
;
using
detail
::
limited_vector
;
namespace
{
using
ivec
=
std
::
vector
<
int
>
;
constexpr
size_t
matrix_size
=
4
;
constexpr
size_t
array_size
=
32
;
constexpr
int
magic_number
=
23
;
constexpr
const
char
*
kernel_name
=
"matrix_square"
;
constexpr
const
char
*
kernel_name_compiler_flag
=
"compiler_flag"
;
constexpr
const
char
*
kernel_name_const
=
"const_mod"
;
constexpr
const
char
*
compiler_flag
=
"-D CAF_OPENCL_TEST_FLAG"
;
constexpr
const
char
*
kernel_source
=
R"__(
__kernel void matrix_square(__global int* matrix,
__global int* 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;
}
)__"
;
constexpr
const
char
*
kernel_source_error
=
R"__(
__kernel void missing(__global int*) {
size_t semicolon
}
)__"
;
constexpr
const
char
*
kernel_source_compiler_flag
=
R"__(
__kernel void compiler_flag(__global int* input,
__global int* output) {
size_t x = get_global_id(0);
# ifdef CAF_OPENCL_TEST_FLAG
output[x] = input[x];
# else
output[x] = 0;
# endif
}
)__"
;
constexpr
const
char
*
kernel_source_const
=
R"__(
__kernel void const_mod(__constant int* input,
__global int* output) {
size_t idx = get_global_id(0);
output[idx] = input[0];
}
)__"
;
}
// namespace <anonymous>
template
<
size_t
Size
>
class
square_matrix
{
public:
static
constexpr
size_t
num_elements
=
Size
*
Size
;
static
void
announce
()
{
caf
::
announce
<
square_matrix
>
(
"square_matrix"
,
&
square_matrix
::
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
);
}
float
&
operator
()(
size_t
column
,
size_t
row
)
{
return
data_
[
column
+
row
*
Size
];
}
const
float
&
operator
()(
size_t
column
,
size_t
row
)
const
{
return
data_
[
column
+
row
*
Size
];
}
typedef
typename
ivec
::
const_iterator
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_
=
std
::
move
(
new_data
);
}
private:
ivec
data_
;
};
template
<
class
T
>
std
::
vector
<
T
>
make_iota_vector
(
size_t
num_elements
)
{
std
::
vector
<
T
>
result
;
result
.
resize
(
num_elements
);
std
::
iota
(
result
.
begin
(),
result
.
end
(),
T
{
0
});
return
result
;
}
template
<
size_t
Size
>
square_matrix
<
Size
>
make_iota_matrix
()
{
square_matrix
<
Size
>
result
;
std
::
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
std
::
string
&
description
,
const
std
::
vector
<
T
>&
expected
,
const
std
::
vector
<
T
>&
result
)
{
auto
cond
=
(
expected
==
result
);
CAF_CHECK
(
cond
);
if
(
!
cond
)
{
CAF_TEST_INFO
(
description
<<
" failed."
);
std
::
cout
<<
"Expected: "
<<
std
::
endl
;
for
(
size_t
i
=
0
;
i
<
expected
.
size
();
++
i
)
{
std
::
cout
<<
" "
<<
expected
[
i
];
}
std
::
cout
<<
std
::
endl
<<
"Received: "
<<
std
::
endl
;
for
(
size_t
i
=
0
;
i
<
result
.
size
();
++
i
)
{
std
::
cout
<<
" "
<<
result
[
i
];
}
std
::
cout
<<
std
::
endl
;
}
}
void
test_opencl_deprecated
()
{
scoped_actor
self
;
const
ivec
expected1
{
56
,
62
,
68
,
74
,
152
,
174
,
196
,
218
,
248
,
286
,
324
,
362
,
344
,
398
,
452
,
506
};
auto
w1
=
spawn_cl
<
ivec
(
ivec
)
>
(
program
::
create
(
kernel_source
),
kernel_name
,
limited_vector
<
size_t
,
3
>
{
matrix_size
,
matrix_size
});
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)"
,
result
,
expected1
);
}
);
auto
w2
=
spawn_cl
<
ivec
(
ivec
)
>
(
kernel_source
,
kernel_name
,
limited_vector
<
size_t
,
3
>
{
matrix_size
,
matrix_size
});
self
->
send
(
w2
,
make_iota_vector
<
int
>
(
matrix_size
*
matrix_size
));
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Simple matrix multiplication using vectors"
,
expected1
,
result
);
}
);
const
matrix_type
expected2
(
std
::
move
(
expected1
));
auto
map_arg
=
[](
message
&
msg
)
->
maybe
<
message
>
{
return
msg
.
apply
(
[](
matrix_type
&
mx
)
{
return
make_message
(
std
::
move
(
mx
.
data
()));
}
);
};
auto
map_res
=
[](
ivec
&
result
)
->
message
{
return
make_message
(
matrix_type
{
std
::
move
(
result
)});
};
auto
w3
=
spawn_cl
<
ivec
(
ivec
)
>
(
program
::
create
(
kernel_source
),
kernel_name
,
map_arg
,
map_res
,
limited_vector
<
size_t
,
3
>
{
matrix_size
,
matrix_size
});
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
());
}
);
auto
w4
=
spawn_cl
<
ivec
(
ivec
)
>
(
kernel_source
,
kernel_name
,
map_arg
,
map_res
,
limited_vector
<
size_t
,
3
>
{
matrix_size
,
matrix_size
});
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
());
}
);
CAF_TEST_INFO
(
"Expecting exception (compiling invalid kernel, "
"semicolon is missing)."
);
try
{
auto
create_error
=
program
::
create
(
kernel_source_error
);
}
catch
(
const
std
::
exception
&
exc
)
{
auto
cond
=
(
strcmp
(
"clBuildProgram: CL_BUILD_PROGRAM_FAILURE"
,
exc
.
what
())
==
0
);
CAF_CHECK
(
cond
);
if
(
!
cond
)
{
CAF_TEST_INFO
(
"Wrong exception cought for program build failure."
);
}
}
// test for opencl compiler flags
auto
prog5
=
program
::
create
(
kernel_source_compiler_flag
,
compiler_flag
);
auto
w5
=
spawn_cl
<
ivec
(
ivec
)
>
(
prog5
,
kernel_name_compiler_flag
,
limited_vector
<
size_t
,
3
>
{
array_size
});
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
);
}
);
// constant memory arguments
const
ivec
arr7
{
magic_number
};
auto
w7
=
spawn_cl
<
ivec
(
ivec
)
>
(
kernel_source_const
,
kernel_name_const
,
limited_vector
<
size_t
,
3
>
{
magic_number
});
self
->
send
(
w7
,
move
(
arr7
));
ivec
expected5
(
magic_number
);
fill
(
begin
(
expected5
),
end
(
expected5
),
magic_number
);
self
->
receive
(
[
&
](
const
ivec
&
result
)
{
check_vector_results
(
"Using const input argument"
,
expected5
,
result
);
}
);
}
CAF_TEST
(
test_opencl_deprecated
)
{
std
::
cout
<<
"Staring deprecated OpenCL test"
<<
std
::
endl
;
announce
<
ivec
>
(
"ivec"
);
matrix_type
::
announce
();
test_opencl_deprecated
();
await_all_actors_done
();
shutdown
();
std
::
cout
<<
"Done with depreacted OpenCL test"
<<
std
::
endl
;
}
CAF_POP_WARNINGS
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