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
e5f44ee4
Commit
e5f44ee4
authored
Jun 15, 2017
by
Joseph Noir
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
Cleanup (addresses review)
parent
f314c210
Changes
21
Hide whitespace changes
Inline
Side-by-side
Showing
21 changed files
with
557 additions
and
572 deletions
+557
-572
libcaf_opencl/caf/opencl/actor_facade.hpp
libcaf_opencl/caf/opencl/actor_facade.hpp
+32
-72
libcaf_opencl/caf/opencl/arguments.hpp
libcaf_opencl/caf/opencl/arguments.hpp
+117
-156
libcaf_opencl/caf/opencl/command.hpp
libcaf_opencl/caf/opencl/command.hpp
+57
-76
libcaf_opencl/caf/opencl/detail/command_helper.hpp
libcaf_opencl/caf/opencl/detail/command_helper.hpp
+63
-0
libcaf_opencl/caf/opencl/detail/core.hpp
libcaf_opencl/caf/opencl/detail/core.hpp
+38
-0
libcaf_opencl/caf/opencl/detail/raw_ptr.hpp
libcaf_opencl/caf/opencl/detail/raw_ptr.hpp
+12
-7
libcaf_opencl/caf/opencl/detail/spawn_helper.hpp
libcaf_opencl/caf/opencl/detail/spawn_helper.hpp
+5
-5
libcaf_opencl/caf/opencl/device.hpp
libcaf_opencl/caf/opencl/device.hpp
+70
-68
libcaf_opencl/caf/opencl/global.hpp
libcaf_opencl/caf/opencl/global.hpp
+1
-1
libcaf_opencl/caf/opencl/manager.hpp
libcaf_opencl/caf/opencl/manager.hpp
+28
-51
libcaf_opencl/caf/opencl/mem_ref.hpp
libcaf_opencl/caf/opencl/mem_ref.hpp
+26
-20
libcaf_opencl/caf/opencl/platform.hpp
libcaf_opencl/caf/opencl/platform.hpp
+10
-10
libcaf_opencl/caf/opencl/program.hpp
libcaf_opencl/caf/opencl/program.hpp
+10
-14
libcaf_opencl/examples/scan.cpp
libcaf_opencl/examples/scan.cpp
+12
-27
libcaf_opencl/src/device.cpp
libcaf_opencl/src/device.cpp
+18
-11
libcaf_opencl/src/global.cpp
libcaf_opencl/src/global.cpp
+1
-1
libcaf_opencl/src/manager.cpp
libcaf_opencl/src/manager.cpp
+13
-12
libcaf_opencl/src/opencl_err.cpp
libcaf_opencl/src/opencl_err.cpp
+1
-1
libcaf_opencl/src/platform.cpp
libcaf_opencl/src/platform.cpp
+6
-4
libcaf_opencl/src/program.cpp
libcaf_opencl/src/program.cpp
+4
-3
libcaf_opencl/test/opencl.cpp
libcaf_opencl/test/opencl.cpp
+33
-33
No files found.
libcaf_opencl/caf/opencl/
opencl_actor
.hpp
→
libcaf_opencl/caf/opencl/
actor_facade
.hpp
View file @
e5f44ee4
...
...
@@ -17,8 +17,8 @@
* http://www.boost.org/LICENSE_1_0.txt. *
******************************************************************************/
#ifndef CAF_OPENCL_
OPENCL_ACTOR
_HPP
#define CAF_OPENCL_
OPENCL_ACTOR
_HPP
#ifndef CAF_OPENCL_
ACTOR_FACADE
_HPP
#define CAF_OPENCL_
ACTOR_FACADE
_HPP
#include <ostream>
#include <iostream>
...
...
@@ -29,69 +29,27 @@
#include "caf/intrusive_ptr.hpp"
#include "caf/detail/int_list.hpp"
#include "caf/detail/type_list.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/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/nd_range.hpp"
#include "caf/opencl/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
#include "caf/opencl/detail/command_helper.hpp"
namespace
caf
{
namespace
opencl
{
class
manager
;
// 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
&
...)
>
;
};
// convert to mem_ref
template
<
class
T
>
struct
to_mem_ref
{
using
type
=
mem_ref
<
T
>
;
};
template
<
class
T
>
struct
to_mem_ref
<
std
::
vector
<
T
>>
{
using
type
=
mem_ref
<
T
>
;
};
template
<
class
T
>
struct
to_mem_ref
<
T
*>
{
using
type
=
mem_ref
<
T
>
;
};
// 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
=
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
...
>
;
};
template
<
bool
PassConfig
,
class
...
Ts
>
class
opencl_actor
:
public
monitorable_actor
{
class
actor_facade
:
public
monitorable_actor
{
public:
using
arg_types
=
detail
::
type_list
<
Ts
...
>
;
using
unpacked_types
=
typename
detail
::
tl_map
<
arg_types
,
extract_type
>::
type
;
...
...
@@ -109,18 +67,18 @@ public:
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
output_function_sig
<
output_types
>::
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
command_sig
<
opencl_actor
,
output_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
<
cl
_mem_ptr
>
;
using
mem_vec
=
std
::
vector
<
detail
::
raw
_mem_ptr
>
;
using
len_vec
=
std
::
vector
<
size_t
>
;
using
out_tup
=
typename
tuple_type_of
<
output_types
>::
type
;
using
out_tup
=
typename
detail
::
tuple_type_of
<
output_types
>::
type
;
const
char
*
name
()
const
override
{
return
"OpenCL actor"
;
...
...
@@ -149,18 +107,18 @@ public:
auto
&
sys
=
actor_conf
.
host
->
system
();
auto
itr
=
prog
->
available_kernels_
.
find
(
kernel_name
);
if
(
itr
==
prog
->
available_kernels_
.
end
())
{
cl
_kernel_ptr
kernel
;
detail
::
raw
_kernel_ptr
kernel
;
kernel
.
reset
(
v2get
(
CAF_CLF
(
clCreateKernel
),
prog
->
program_
.
get
(),
kernel_name
),
false
);
return
make_actor
<
opencl_actor
,
actor
>
(
sys
.
next_actor_id
(),
sys
.
node
(),
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
<
opencl_actor
,
actor
>
(
sys
.
next_actor_id
(),
sys
.
node
(),
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
),
...
...
@@ -219,11 +177,12 @@ public:
void
enqueue
(
strong_actor_ptr
sender
,
message_id
mid
,
message
content
,
execution_unit
*
host
)
override
{
CAF_LOG_TRACE
(
""
);
enqueue
(
make_mailbox_element
(
sender
,
mid
,
{},
std
::
move
(
content
)),
host
);
enqueue
(
make_mailbox_element
(
std
::
move
(
sender
),
mid
,
{},
std
::
move
(
content
)),
host
);
}
opencl_actor
(
actor_config
actor_conf
,
const
program_ptr
prog
,
cl
_kernel_ptr
kernel
,
nd_range
range
,
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
)
:
monitorable_actor
(
actor_conf
),
...
...
@@ -243,7 +202,8 @@ public:
}
void
add_kernel_arguments
(
evnt_vec
&
,
mem_vec
&
,
mem_vec
&
,
mem_vec
&
,
out_tup
&
,
len_vec
&
,
message
&
,
detail
::
int_list
<>
)
{
out_tup
&
,
len_vec
&
,
message
&
,
detail
::
int_list
<>
)
{
// nop
}
...
...
@@ -254,7 +214,7 @@ public:
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
caf
::
detail
::
tl_at
<
processing_list
,
I
>::
type
;
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
...
...
@@ -340,7 +300,7 @@ public:
sizeof
(
cl_mem
),
static_cast
<
const
void
*>
(
&
buffer
));
events
.
push_back
(
event
);
std
::
get
<
OutPos
>
(
result
)
=
mem_ref
<
value_type
>
{
len
,
queue_
,
cl
_mem_ptr
{
buffer
,
false
},
len
,
queue_
,
detail
::
raw
_mem_ptr
{
buffer
,
false
},
size_t
{
CL_MEM_READ_WRITE
|
CL_MEM_HOST_READ_ONLY
},
nullptr
};
}
...
...
@@ -435,7 +395,7 @@ public:
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
;
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
),
...
...
@@ -468,12 +428,12 @@ public:
template
<
class
Fun
>
size_t
argument_length
(
Fun
&
f
,
message
&
m
,
size_t
fallback
)
{
auto
length
=
f
(
m
);
return
length
&&
(
*
length
>
0
)
?
*
length
:
fallback
;
return
length
&&
(
*
length
>
0
)
?
*
length
:
fallback
;
}
// Map function requires only the message as argument
template
<
bool
Q
=
PassConfig
>
typename
std
::
enable_if
<!
Q
,
bool
>::
type
map_arguments
(
message
&
content
)
{
detail
::
enable_if_t
<!
Q
,
bool
>
map_arguments
(
message
&
content
)
{
if
(
map_args_
)
{
auto
mapped
=
map_args_
(
content
);
if
(
!
mapped
)
{
...
...
@@ -487,7 +447,7 @@ public:
// Map function requires reference to config as well as the message
template
<
bool
Q
=
PassConfig
>
typename
std
::
enable_if
<
Q
,
bool
>::
type
map_arguments
(
message
&
content
)
{
detail
::
enable_if_t
<
Q
,
bool
>
map_arguments
(
message
&
content
)
{
if
(
map_args_
)
{
auto
mapped
=
map_args_
(
range_
,
content
);
if
(
!
mapped
)
{
...
...
@@ -499,10 +459,10 @@ public:
return
true
;
}
cl
_kernel_ptr
kernel_
;
cl
_program_ptr
program_
;
cl
_context_ptr
context_
;
cl
_command_queue_ptr
queue_
;
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_
;
...
...
@@ -512,5 +472,5 @@ public:
}
// namespace opencl
}
// namespace caf
#endif // CAF_OPENCL_OPENCL_ACTOR_HPP
#endif // CAF_OPENCL_ACTOR_FACADE_HPP
libcaf_opencl/caf/opencl/arguments.hpp
View file @
e5f44ee4
...
...
@@ -28,22 +28,50 @@
#include "caf/optional.hpp"
#include "caf/opencl/mem_ref.hpp"
#include "caf/opencl/detail/core.hpp"
namespace
caf
{
namespace
opencl
{
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
// 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
{
};
/// 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
{
};
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
{
};
struct
hidden
{};
/// Use as a default way to calculate output size. 0 will be set to the number
/// of work items at runtime.
...
...
@@ -54,146 +82,128 @@ struct dummy_size_calculator {
}
};
/// Mark an a spawn template argument as input only
/// 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
{
static_assert
(
std
::
is_same
<
Tag
,
val
>::
value
||
std
::
is_same
<
Tag
,
mref
>::
value
,
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
=
typename
std
::
decay
<
Arg
>::
type
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
};
/// Mark a
n a spawn template
argument as input and output
/// Mark a
spawn
argument as input and output
template
<
class
Arg
,
class
TagIn
=
val
,
class
TagOut
=
val
>
struct
in_out
{
struct
in_out
:
arg_tag
,
input_tag
,
output_tag
{
static_assert
(
std
::
is_same
<
TagIn
,
val
>::
value
||
std
::
is_same
<
TagIn
,
mref
>::
value
,
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
,
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
=
typename
std
::
decay
<
Arg
>::
type
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
};
/// Mark a spawn argument as output only
template
<
class
Arg
,
class
Tag
=
val
>
struct
out
{
static_assert
(
std
::
is_same
<
Tag
,
val
>::
value
||
std
::
is_same
<
Tag
,
mref
>::
value
,
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
=
typename
std
::
decay
<
Arg
>::
type
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
out
()
=
default
;
template
<
class
F
>
out
(
F
fun
)
{
fun_
=
[
fun
](
message
&
msg
)
->
optional
<
size_t
>
{
auto
res
=
msg
.
apply
(
fun
);
size_t
result
;
if
(
res
)
{
res
->
apply
([
&
](
size_t
x
)
{
result
=
x
;
});
return
result
;
}
return
none
;
};
out
(
F
fun
)
:
fun_
{
detail
::
res_or_none
<
size_t
>
(
fun
)}
{
// nop
}
optional
<
size_t
>
operator
()(
message
&
msg
)
const
{
return
fun_
?
fun_
(
msg
)
:
0UL
;
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
{
using
arg_type
=
typename
std
::
decay
<
Arg
>::
type
;
struct
scratch
:
arg_tag
,
requires_size_tag
{
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
scratch
()
=
default
;
template
<
class
F
>
scratch
(
F
fun
)
{
fun_
=
[
fun
](
message
&
msg
)
->
optional
<
size_t
>
{
auto
res
=
msg
.
apply
(
fun
);
size_t
result
;
if
(
res
)
{
res
->
apply
([
&
](
size_t
x
)
{
result
=
x
;
});
return
result
;
}
return
none
;
};
scratch
(
F
fun
)
:
fun_
{
detail
::
res_or_none
<
size_t
>
(
fun
)}
{
// nop
}
optional
<
size_t
>
operator
()(
message
&
msg
)
const
{
return
fun_
?
fun_
(
msg
)
:
0UL
;
return
detail
::
try_apply_fun
(
fun_
,
msg
,
0UL
)
;
}
std
::
function
<
optional
<
size_t
>
(
message
&
)
>
fun_
;
};
/// Argument placed in local memory. Cannot be initalized from the CPU, but
/// requires a size that is calculated depending on the input.
/// 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
{
using
arg_type
=
typename
std
::
decay
<
Arg
>::
type
;
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_
=
[
fun
](
message
&
msg
)
->
optional
<
size_t
>
{
auto
res
=
msg
.
apply
(
fun
);
size_t
result
;
if
(
res
)
{
res
->
apply
([
&
](
size_t
x
)
{
result
=
x
;
});
return
result
;
}
return
none
;
};
local
(
size_t
size
,
F
fun
)
:
size_
(
size
),
fun_
{
detail
::
res_or_none
<
size_t
>
(
fun
)}
{
// nop
}
local
(
size_t
size
)
:
size_
(
size
)
{
}
size_t
operator
()(
message
&
msg
)
const
{
if
(
fun_
)
{
auto
res
=
fun_
(
msg
);
if
(
res
)
return
*
res
;
}
return
size_
;
return
detail
::
try_apply_fun
(
fun_
,
msg
,
size_
);
}
size_t
size_
;
std
::
function
<
optional
<
size_t
>
(
message
&
)
>
fun_
;
};
///
Argument placed in private memory. Requires a default value but can
///
alternatively be calculated depending on the input through a function
///
passed to the constructor
.
///
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
{
static_assert
(
std
::
is_same
<
Tag
,
val
>::
value
||
std
::
is_same
<
Tag
,
hidden
>::
value
,
"Argument of type `priv` must be returned as value or 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
=
typename
std
::
decay
<
Arg
>::
type
;
using
arg_type
=
detail
::
decay_t
<
Arg
>
;
priv
()
=
default
;
template
<
class
F
>
priv
(
Arg
val
,
F
fun
)
:
value_
(
val
)
{
static_assert
(
std
::
is_same
<
Tag
,
hidden
>::
value
,
"Argument of type `priv` can only be initialized with a value"
"if it is manged by the actor, i.e., tagged as hidden."
);
fun_
=
[
fun
](
message
&
msg
)
->
optional
<
Arg
>
{
auto
res
=
msg
.
apply
(
fun
);
Arg
result
;
if
(
res
)
{
res
->
apply
([
&
](
Arg
x
)
{
result
=
x
;
});
return
result
;
}
return
none
;
};
}
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 manged by the actor, i.e., tagged as hidden."
);
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
{
if
(
fun_
)
{
auto
res
=
fun_
(
msg
);
if
(
res
)
return
*
res
;
}
return
value_
;
return
detail
::
try_apply_fun
(
fun_
,
msg
,
value_
);
}
Arg
value_
;
std
::
function
<
optional
<
Arg
>
(
message
&
)
>
fun_
;
...
...
@@ -212,77 +222,27 @@ struct carr_to_vec<T*> {
/// Filter types for any argument type.
template
<
class
T
>
struct
is_opencl_arg
:
std
::
false_type
{};
template
<
class
T
,
class
Tag
>
struct
is_opencl_arg
<
in
<
T
,
Tag
>>
:
std
::
true_type
{};
template
<
class
T
,
class
TagIn
,
class
TagOut
>
struct
is_opencl_arg
<
in_out
<
T
,
TagIn
,
TagOut
>>
:
std
::
true_type
{};
template
<
class
T
,
class
Tag
>
struct
is_opencl_arg
<
out
<
T
,
Tag
>>
:
std
::
true_type
{};
template
<
class
T
>
struct
is_opencl_arg
<
scratch
<
T
>>
:
std
::
true_type
{};
template
<
class
T
>
struct
is_opencl_arg
<
local
<
T
>>
:
std
::
true_type
{};
template
<
class
T
,
class
Tag
>
struct
is_opencl_arg
<
priv
<
T
,
Tag
>>
:
std
::
true_type
{};
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
::
false_type
{};
template
<
class
T
,
class
Tag
>
struct
is_input_arg
<
in
<
T
,
Tag
>>
:
std
::
true_type
{};
template
<
class
T
,
class
TagIn
,
class
TagOut
>
struct
is_input_arg
<
in_out
<
T
,
TagIn
,
TagOut
>>
:
std
::
true_type
{};
template
<
class
T
>
struct
is_input_arg
<
priv
<
T
,
val
>>
:
std
::
true_type
{};
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
::
false_type
{};
template
<
class
T
,
class
Tag
>
struct
is_output_arg
<
out
<
T
,
Tag
>>
:
std
::
true_type
{};
template
<
class
T
,
class
TagIn
,
class
TagOut
>
struct
is_output_arg
<
in_out
<
T
,
TagIn
,
TagOut
>>
:
std
::
true_type
{};
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
::
false_type
{};
template
<
class
T
,
class
Tag
>
struct
requires_size_arg
<
out
<
T
,
Tag
>>
:
std
::
true_type
{};
template
<
class
T
>
struct
requires_size_arg
<
scratch
<
T
>>
:
std
::
true_type
{};
template
<
class
T
>
struct
requires_size_arg
<
local
<
T
>>
:
std
::
true_type
{};
//template <class T, class Tag>
//struct requires_size_arg<priv<T, Tag>> : std::true_type {};
struct
requires_size_arg
:
std
::
is_base_of
<
requires_size_tag
,
T
>
{};
/// Filter mem_refs
template
<
class
T
>
struct
is_ref_type
:
std
::
false_type
{};
template
<
class
T
>
struct
is_ref_type
<
mem_ref
<
T
>>
:
std
::
true_type
{};
struct
is_ref_type
:
std
::
is_base_of
<
is_ref_tag
,
T
>
{};
template
<
class
T
>
struct
is_val_type
:
std
::
true_type
{};
template
<
class
T
>
struct
is_val_type
<
mem_ref
<
T
>>
:
std
::
false_type
{};
struct
is_val_type
:
std
::
integral_constant
<
bool
,
!
std
::
is_base_of
<
is_ref_tag
,
T
>::
value
>
{};
/// extract types
template
<
class
T
>
...
...
@@ -290,32 +250,32 @@ struct extract_type { };
template
<
class
T
,
class
Tag
>
struct
extract_type
<
in
<
T
,
Tag
>>
{
using
type
=
typename
std
::
decay
<
typename
carr_to_vec
<
T
>::
type
>::
type
;
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
=
typename
std
::
decay
<
typename
carr_to_vec
<
T
>::
type
>::
type
;
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
,
class
Tag
>
struct
extract_type
<
out
<
T
,
Tag
>>
{
using
type
=
typename
std
::
decay
<
typename
carr_to_vec
<
T
>::
type
>::
type
;
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
>
struct
extract_type
<
scratch
<
T
>>
{
using
type
=
typename
std
::
decay
<
typename
carr_to_vec
<
T
>::
type
>::
type
;
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
>
struct
extract_type
<
local
<
T
>>
{
using
type
=
typename
std
::
decay
<
typename
carr_to_vec
<
T
>::
type
>::
type
;
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
template
<
class
T
,
class
Tag
>
struct
extract_type
<
priv
<
T
,
Tag
>>
{
using
type
=
typename
std
::
decay
<
typename
carr_to_vec
<
T
>::
type
>::
type
;
using
type
=
detail
::
decay_t
<
typename
carr_to_vec
<
T
>::
type
>
;
};
/// extract type expected in an incoming message
...
...
@@ -504,7 +464,8 @@ struct cl_arg_info_list_impl<detail::type_list<Args...>,
};
template
<
class
...
Args
,
int
InCounter
,
int
OutCounter
>
struct
cl_arg_info_list_impl
<
detail
::
type_list
<
Args
...
>
,
detail
::
type_list
<>
,
struct
cl_arg_info_list_impl
<
detail
::
type_list
<
Args
...
>
,
detail
::
type_list
<>
,
InCounter
,
OutCounter
>
{
using
type
=
detail
::
type_list
<
Args
...
>
;
};
...
...
libcaf_opencl/caf/opencl/command.hpp
View file @
e5f44ee4
...
...
@@ -34,10 +34,12 @@
#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/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/nd_range.hpp"
#include "caf/opencl/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace
caf
{
namespace
opencl
{
...
...
@@ -54,9 +56,9 @@ public:
command
(
response_promise
promise
,
strong_actor_ptr
parent
,
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
cl
_mem_ptr
>
inputs
,
std
::
vector
<
cl
_mem_ptr
>
outputs
,
std
::
vector
<
cl
_mem_ptr
>
scratches
,
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
,
...
...
@@ -68,7 +70,7 @@ public:
input_buffers_
(
std
::
move
(
inputs
)),
output_buffers_
(
std
::
move
(
outputs
)),
scratch_buffers_
(
std
::
move
(
scratches
)),
results_
(
output_tuple
),
results_
(
std
::
move
(
output_tuple
)
),
msg_
(
std
::
move
(
msg
)),
range_
(
std
::
move
(
range
))
{
// nop
...
...
@@ -85,19 +87,16 @@ public:
}
}
/// Enqueue the kernel for execution, schedule reading of the results and
/// 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
>
typename
std
::
enable_if
<
!
detail
::
tl_forall
<
Q
,
is_ref_type
>::
value
,
void
>::
type
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
(
"
command::enqueue() mixed
"
);
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
();
...
...
@@ -105,8 +104,8 @@ public:
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
();
cl_int
err
=
clEnqueueNDRangeKerne
l
(
parent
->
queue_
.
get
(),
parent
->
kernel_
.
get
(),
auto
success
=
invoke_c
l
(
clEnqueueNDRangeKernel
,
parent
->
queue_
.
get
(),
parent
->
kernel_
.
get
(),
static_cast
<
unsigned
int
>
(
range_
.
dimensions
().
size
()),
data_or_nullptr
(
range_
.
offsets
()),
data_or_nullptr
(
range_
.
dimensions
()),
...
...
@@ -115,51 +114,34 @@ public:
(
mem_in_events_
.
empty
()
?
nullptr
:
mem_in_events_
.
data
()),
&
mem_out_events_
.
back
()
);
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG_ERROR
(
"clEnqueueNDRangeKernel: "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
this
->
deref
();
if
(
!
success
)
return
;
}
size_t
pos
=
0
;
CAF_ASSERT
(
!
mem_out_events_
.
empty
());
enqueue_read_buffers
(
pos
,
mem_out_events_
,
detail
::
get_indices
(
results_
));
enqueue_read_buffers
(
pos
,
mem_out_events_
,
detail
::
get_indices
(
results_
));
CAF_ASSERT
(
mem_out_events_
.
size
()
>
1
);
cl_event
marker_event
;
#if defined(__APPLE__)
err
=
clEnqueueMarkerWithWaitList
(
parent
->
queue_
.
get
(),
static_cast
<
unsigned
int
>
(
mem_out_events_
.
size
()),
mem_out_events_
.
data
(),
&
marker_event
);
std
::
string
name
=
"clEnqueueMarkerWithWaitList"
;
success
=
invoke_cl
(
clEnqueueMarkerWithWaitList
,
parent
->
queue_
.
get
(),
static_cast
<
unsigned
int
>
(
mem_out_events_
.
size
()),
mem_out_events_
.
data
(),
&
marker_event
);
#else
err
=
clEnqueueMarker
(
parent
->
queue_
.
get
(),
&
marker_event
);
std
::
string
name
=
"clEnqueueMarker"
;
success
=
invoke_cl
(
clEnqueueMarker
,
parent
->
queue_
.
get
(),
&
marker_event
);
#endif
callback_
.
reset
(
marker_event
,
false
);
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG_ERROR
(
name
<<
": "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
this
->
deref
();
// callback is not set
if
(
!
success
)
return
;
}
err
=
clSetEventCallback
(
callback_
.
get
(),
CL_COMPLETE
,
[](
cl_event
,
cl_int
,
void
*
data
)
{
auto
cmd
=
reinterpret_cast
<
command
*>
(
data
);
cmd
->
handle_results
();
cmd
->
deref
();
},
this
);
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG_ERROR
(
"clSetEventCallback: "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
this
->
deref
();
// callback is not set
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
;
}
err
=
clFlush
(
parent
->
queue_
.
get
());
if
(
err
!=
CL_SUCCESS
)
{
if
(
clFlush
(
parent
->
queue_
.
get
())
!=
CL_SUCCESS
)
CAF_LOG_ERROR
(
"clFlush: "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
}
}
/// Enqueue the kernel for execution and send the mem_refs relating to the
...
...
@@ -167,22 +149,19 @@ public:
/// once the execution is finished. Only called if the results only consist
/// of mem_ref types.
template
<
class
Q
=
result_types
>
typename
std
::
enable_if
<
detail
::
tl_forall
<
Q
,
is_ref_type
>::
value
,
void
>::
type
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
(
"
command::enqueue() all references
"
);
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
;
cl_int
err
=
clEnqueueNDRangeKerne
l
(
parent
->
queue_
.
get
(),
parent
->
kernel_
.
get
(),
auto
success
=
invoke_c
l
(
clEnqueueNDRangeKernel
,
parent
->
queue_
.
get
(),
parent
->
kernel_
.
get
(),
static_cast
<
cl_uint
>
(
range_
.
dimensions
().
size
()),
data_or_nullptr
(
range_
.
offsets
()),
data_or_nullptr
(
range_
.
dimensions
()),
...
...
@@ -192,25 +171,16 @@ public:
&
execution_event
);
callback_
.
reset
(
execution_event
,
false
);
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG_ERROR
(
"clEnqueueNDRangeKernel: "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
this
->
deref
();
if
(
!
success
)
return
;
}
err
=
clSetEventCallback
(
callback_
.
get
(),
CL_COMPLETE
,
[](
cl_event
,
cl_int
,
void
*
data
)
{
auto
c
=
reinterpret_cast
<
command
*>
(
data
);
c
->
deref
();
},
this
);
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOG_ERROR
(
"clSetEventCallback: "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
this
->
deref
();
// callback is not set
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
;
}
err
=
clFlush
(
parent
->
queue_
.
get
());
if
(
err
!=
CL_SUCCESS
)
if
(
clFlush
(
parent
->
queue_
.
get
())
!=
CL_SUCCESS
)
CAF_LOG_ERROR
(
"clFlush: "
<<
CAF_ARG
(
get_opencl_error
(
err
)));
auto
msg
=
msg_adding_event
{
callback_
}(
results_
);
promise_
.
deliver
(
std
::
move
(
msg
));
...
...
@@ -231,7 +201,7 @@ private:
events
.
data
(),
&
events
.
back
());
if
(
err
!=
CL_SUCCESS
)
{
this
->
deref
();
// failed to enqueue command
throw
std
::
runtime_error
(
"clEnqueueReadBuffer: "
+
get_
opencl_error
(
err
));
throw
std
::
runtime_error
(
"clEnqueueReadBuffer: "
+
opencl_error
(
err
));
}
pos
+=
1
;
}
...
...
@@ -263,15 +233,26 @@ private:
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_
;
cl
_event_ptr
callback_
;
std
::
vector
<
cl
_mem_ptr
>
input_buffers_
;
std
::
vector
<
cl
_mem_ptr
>
output_buffers_
;
std
::
vector
<
cl
_mem_ptr
>
scratch_buffers_
;
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_
;
...
...
libcaf_opencl/caf/opencl/detail/command_helper.hpp
0 → 100644
View file @
e5f44ee4
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* 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_DETAIL_COMMAND_HELPER_HPP
#define CAF_OPENCL_DETAIL_COMMAND_HELPER_HPP
#include "caf/detail/type_list.hpp"
#include "caf/opencl/arguments.hpp"
namespace
caf
{
namespace
opencl
{
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
=
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 opencl
}
// namespace caf
#endif // CAF_OPENCL_DETAIL_COMMAND_HELPER_HPP
libcaf_opencl/caf/opencl/detail/core.hpp
0 → 100644
View file @
e5f44ee4
/******************************************************************************
* ____ _ _____ *
* / ___| / \ | ___| C++ *
* | | / _ \ | |_ Actor *
* | |___ / ___ \| _| Framework *
* \____/_/ \_|_| *
* *
* Copyright (C) 2011 - 2016 *
* 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_CORE_HPP
#define CAF_OPENCL_CORE_HPP
namespace
caf
{
namespace
detail
{
}
// namespace detail
namespace
opencl
{
namespace
detail
{
using
namespace
caf
::
detail
;
}
// namespace detail
}
// namespace opencl
}
// namespace caf
#endif // CAF_OPENCL_CORE_HPP
libcaf_opencl/caf/opencl/
smart
_ptr.hpp
→
libcaf_opencl/caf/opencl/
detail/raw
_ptr.hpp
View file @
e5f44ee4
...
...
@@ -33,24 +33,29 @@
inline void intrusive_ptr_release(cltype ptr) { clrelease(ptr); } \
namespace caf { \
namespace opencl { \
namespace detail { \
using aliasname = intrusive_ptr<std::remove_pointer<cltype>::type>; \
}
/* namespace detail */
\
}
/* namespace opencl */
\
} // namespace caf
CAF_OPENCL_PTR_ALIAS
(
cl_mem_ptr
,
cl_mem
,
clRetainMemObject
,
clReleaseMemObject
)
CAF_OPENCL_PTR_ALIAS
(
cl_event_ptr
,
cl_event
,
clRetainEvent
,
clReleaseEven
t
)
CAF_OPENCL_PTR_ALIAS
(
raw_mem_ptr
,
cl_mem
,
clRetainMemObject
,
clReleaseMemObjec
t
)
CAF_OPENCL_PTR_ALIAS
(
cl_kernel_ptr
,
cl_kernel
,
clRetainKernel
,
clReleaseKernel
)
CAF_OPENCL_PTR_ALIAS
(
raw_event_ptr
,
cl_event
,
clRetainEvent
,
clReleaseEvent
)
CAF_OPENCL_PTR_ALIAS
(
cl_context_ptr
,
cl_context
,
clRetainContext
,
clReleaseContext
)
CAF_OPENCL_PTR_ALIAS
(
raw_kernel_ptr
,
cl_kernel
,
clRetainKernel
,
clReleaseKernel
)
CAF_OPENCL_PTR_ALIAS
(
cl_program_ptr
,
cl_program
,
clRetainProgram
,
clReleaseProgram
)
CAF_OPENCL_PTR_ALIAS
(
raw_context_ptr
,
cl_context
,
clRetainContext
,
clReleaseContext
)
CAF_OPENCL_PTR_ALIAS
(
cl_device_ptr
,
cl_device_id
,
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
(
cl
_command_queue_ptr
,
cl_command_queue
,
CAF_OPENCL_PTR_ALIAS
(
raw
_command_queue_ptr
,
cl_command_queue
,
clRetainCommandQueue
,
clReleaseCommandQueue
)
#endif // CAF_OPENCL_SMART_PTR_HPP
libcaf_opencl/caf/opencl/detail/spawn_helper.hpp
View file @
e5f44ee4
...
...
@@ -20,7 +20,7 @@
#ifndef CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
#define CAF_OPENCL_DETAIL_SPAWN_HELPER_HPP
#include "caf/opencl/
opencl_actor
.hpp"
#include "caf/opencl/
actor_facade
.hpp"
namespace
caf
{
namespace
opencl
{
...
...
@@ -30,7 +30,7 @@ struct tuple_construct { };
template
<
bool
PassConfig
,
class
...
Ts
>
struct
cl_spawn_helper
{
using
impl
=
opencl
::
opencl_actor
<
PassConfig
,
Ts
...
>
;
using
impl
=
opencl
::
actor_facade
<
PassConfig
,
Ts
...
>
;
using
map_in_fun
=
typename
impl
::
input_mapping
;
using
map_out_fun
=
typename
impl
::
output_mapping
;
...
...
@@ -40,7 +40,7 @@ struct cl_spawn_helper {
return
actor_cast
<
actor
>
(
impl
::
create
(
std
::
move
(
actor_cfg
),
p
,
fn
,
range
,
map_in_fun
{},
map_out_fun
{},
std
::
move
(
xs
)...));
std
::
forward
<
Ts
>
(
xs
)...));
}
actor
operator
()(
actor_config
actor_cfg
,
const
opencl
::
program_ptr
p
,
const
char
*
fn
,
const
opencl
::
nd_range
&
range
,
...
...
@@ -49,7 +49,7 @@ struct cl_spawn_helper {
p
,
fn
,
range
,
std
::
move
(
map_input
),
map_out_fun
{},
std
::
move
(
xs
)...));
std
::
forward
<
Ts
>
(
xs
)...));
}
actor
operator
()(
actor_config
actor_cfg
,
const
opencl
::
program_ptr
p
,
const
char
*
fn
,
const
opencl
::
nd_range
&
range
,
...
...
@@ -59,7 +59,7 @@ struct cl_spawn_helper {
p
,
fn
,
range
,
std
::
move
(
map_input
),
std
::
move
(
map_output
),
std
::
move
(
xs
)...));
std
::
forward
<
Ts
>
(
xs
)...));
}
};
...
...
libcaf_opencl/caf/opencl/device.hpp
View file @
e5f44ee4
...
...
@@ -25,9 +25,10 @@
#include "caf/sec.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace
caf
{
namespace
opencl
{
...
...
@@ -58,10 +59,11 @@ public:
size_t
buffer_size
=
sizeof
(
T
)
*
num_elements
;
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
flags
,
buffer_size
,
nullptr
);
cl_event_ptr
event
{
v1get
<
cl_event
>
(
CAF_CLF
(
clEnqueueWriteBuffer
),
queue_
.
get
(),
buffer
,
blocking
,
cl_uint
{
0
},
buffer_size
,
data
.
data
()),
false
};
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
)};
}
...
...
@@ -92,7 +94,7 @@ public:
buffer_size
,
prev_events
.
size
(),
prev_events
.
data
(),
&
event
);
if
(
err
!=
CL_SUCCESS
)
return
make_error
(
sec
::
runtime_error
,
get_
opencl_error
(
err
));
return
make_error
(
sec
::
runtime_error
,
opencl_error
(
err
));
// callback to release the previous event
if
(
e
)
{
err
=
clSetEventCallback
(
event
,
CL_COMPLETE
,
...
...
@@ -103,7 +105,7 @@ public:
},
e
);
if
(
err
!=
CL_SUCCESS
)
return
make_error
(
sec
::
runtime_error
,
get_
opencl_error
(
err
));
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
),
...
...
@@ -111,80 +113,80 @@ public:
}
/// Initialize a new device in a context using a specific device_id
static
device_ptr
create
(
const
cl
_context_ptr
&
context
,
const
cl
_device_ptr
&
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
get_
id
()
const
;
inline
unsigned
id
()
const
;
/// Returns device info on CL_DEVICE_ADDRESS_BITS
inline
cl_uint
get_
address_bits
()
const
;
inline
cl_uint
address_bits
()
const
;
/// Returns device info on CL_DEVICE_ENDIAN_LITTLE
inline
cl_bool
get_
little_endian
()
const
;
inline
cl_bool
little_endian
()
const
;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHE_SIZE
inline
cl_ulong
g
et_g
lobal_mem_cache_size
()
const
;
inline
cl_ulong
global_mem_cache_size
()
const
;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE
inline
cl_uint
g
et_g
lobal_mem_cacheline_size
()
const
;
inline
cl_uint
global_mem_cacheline_size
()
const
;
/// Returns device info on CL_DEVICE_GLOBAL_MEM_SIZE
inline
cl_ulong
g
et_g
lobal_mem_size
()
const
;
inline
cl_ulong
global_mem_size
()
const
;
/// Returns device info on CL_DEVICE_HOST_UNIFIED_MEMORY
inline
cl_bool
get_
host_unified_memory
()
const
;
inline
cl_bool
host_unified_memory
()
const
;
/// Returns device info on CL_DEVICE_LOCAL_MEM_SIZE
inline
cl_ulong
get_
local_mem_size
()
const
;
inline
cl_ulong
local_mem_size
()
const
;
/// Returns device info on CL_DEVICE_LOCAL_MEM_TYPE
inline
cl_ulong
get_
local_mem_type
()
const
;
inline
cl_ulong
local_mem_type
()
const
;
/// Returns device info on CL_DEVICE_MAX_CLOCK_FREQUENCY
inline
cl_uint
get_
max_clock_frequency
()
const
;
inline
cl_uint
max_clock_frequency
()
const
;
/// Returns device info on CL_DEVICE_MAX_COMPUTE_UNITS
inline
cl_uint
get_
max_compute_units
()
const
;
inline
cl_uint
max_compute_units
()
const
;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_ARGS
inline
cl_uint
get_
max_constant_args
()
const
;
inline
cl_uint
max_constant_args
()
const
;
/// Returns device info on CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE
inline
cl_ulong
get_
max_constant_buffer_size
()
const
;
inline
cl_ulong
max_constant_buffer_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_MEM_ALLOC_SIZE
inline
cl_ulong
get_
max_mem_alloc_size
()
const
;
inline
cl_ulong
max_mem_alloc_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_PARAMETER_SIZE
inline
size_t
get_
max_parameter_size
()
const
;
inline
size_t
max_parameter_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_WORK_GROUP_SIZE
inline
size_t
get_
max_work_group_size
()
const
;
inline
size_t
max_work_group_size
()
const
;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
inline
cl_uint
get_
max_work_item_dimensions
()
const
;
inline
cl_uint
max_work_item_dimensions
()
const
;
/// Returns device info on CL_DEVICE_PROFILING_TIMER_RESOLUTION
inline
size_t
get_
profiling_timer_resolution
()
const
;
inline
size_t
profiling_timer_resolution
()
const
;
/// Returns device info on CL_DEVICE_MAX_WORK_ITEM_SIZES
inline
const
dim_vec
&
get_
max_work_item_sizes
()
const
;
inline
const
dim_vec
&
max_work_item_sizes
()
const
;
/// Returns device info on CL_DEVICE_TYPE
inline
device_type
get_device_
type
()
const
;
inline
device_type
type
()
const
;
/// Returns device info on CL_DEVICE_EXTENSIONS
inline
const
std
::
vector
<
std
::
string
>&
get_
extensions
()
const
;
inline
const
std
::
vector
<
std
::
string
>&
extensions
()
const
;
/// Returns device info on CL_DEVICE_OPENCL_C_VERSION
inline
const
std
::
string
&
get_
opencl_c_version
()
const
;
inline
const
std
::
string
&
opencl_c_version
()
const
;
/// Returns device info on CL_DEVICE_VENDOR
inline
const
std
::
string
&
get_
device_vendor
()
const
;
inline
const
std
::
string
&
device_vendor
()
const
;
/// Returns device info on CL_DEVICE_VERSION
inline
const
std
::
string
&
get_
device_version
()
const
;
inline
const
std
::
string
&
device_version
()
const
;
/// Returns device info on CL_DRIVER_VERSION
inline
const
std
::
string
&
get_
driver_version
()
const
;
inline
const
std
::
string
&
driver_version
()
const
;
/// Returns device info on CL_DEVICE_NAME
inline
const
std
::
string
&
get_
name
()
const
;
inline
const
std
::
string
&
name
()
const
;
private:
device
(
cl_device_ptr
device_id
,
cl
_command_queue_ptr
queue
,
cl
_context_ptr
context
,
unsigned
id
);
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
cl
_device_ptr
&
device_id
,
unsigned
info_flag
)
{
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
cl
_device_ptr
&
device_id
,
static
std
::
string
info_string
(
const
detail
::
raw
_device_ptr
&
device_id
,
unsigned
info_flag
);
cl
_device_ptr
device_id_
;
cl
_command_queue_ptr
queue_
;
cl
_context_ptr
context_
;
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
...
...
@@ -221,107 +223,107 @@ private:
* implementation of inline member functions *
\******************************************************************************/
inline
unsigned
device
::
get_
id
()
const
{
inline
unsigned
device
::
id
()
const
{
return
id_
;
}
inline
cl_uint
device
::
get_
address_bits
()
const
{
inline
cl_uint
device
::
address_bits
()
const
{
return
address_bits_
;
}
inline
cl_bool
device
::
get_
little_endian
()
const
{
inline
cl_bool
device
::
little_endian
()
const
{
return
little_endian_
;
}
inline
cl_ulong
device
::
g
et_g
lobal_mem_cache_size
()
const
{
inline
cl_ulong
device
::
global_mem_cache_size
()
const
{
return
global_mem_cache_size_
;
}
inline
cl_uint
device
::
g
et_g
lobal_mem_cacheline_size
()
const
{
inline
cl_uint
device
::
global_mem_cacheline_size
()
const
{
return
global_mem_cacheline_size_
;
}
inline
cl_ulong
device
::
g
et_g
lobal_mem_size
()
const
{
inline
cl_ulong
device
::
global_mem_size
()
const
{
return
global_mem_size_
;
}
inline
cl_bool
device
::
get_
host_unified_memory
()
const
{
inline
cl_bool
device
::
host_unified_memory
()
const
{
return
host_unified_memory_
;
}
inline
cl_ulong
device
::
get_
local_mem_size
()
const
{
inline
cl_ulong
device
::
local_mem_size
()
const
{
return
local_mem_size_
;
}
inline
cl_ulong
device
::
get_
local_mem_type
()
const
{
inline
cl_ulong
device
::
local_mem_type
()
const
{
return
local_mem_size_
;
}
inline
cl_uint
device
::
get_
max_clock_frequency
()
const
{
inline
cl_uint
device
::
max_clock_frequency
()
const
{
return
max_clock_frequency_
;
}
inline
cl_uint
device
::
get_
max_compute_units
()
const
{
inline
cl_uint
device
::
max_compute_units
()
const
{
return
max_compute_units_
;
}
inline
cl_uint
device
::
get_
max_constant_args
()
const
{
inline
cl_uint
device
::
max_constant_args
()
const
{
return
max_constant_args_
;
}
inline
cl_ulong
device
::
get_
max_constant_buffer_size
()
const
{
inline
cl_ulong
device
::
max_constant_buffer_size
()
const
{
return
max_constant_buffer_size_
;
}
inline
cl_ulong
device
::
get_
max_mem_alloc_size
()
const
{
inline
cl_ulong
device
::
max_mem_alloc_size
()
const
{
return
max_mem_alloc_size_
;
}
inline
size_t
device
::
get_
max_parameter_size
()
const
{
inline
size_t
device
::
max_parameter_size
()
const
{
return
max_parameter_size_
;
}
inline
size_t
device
::
get_
max_work_group_size
()
const
{
inline
size_t
device
::
max_work_group_size
()
const
{
return
max_work_group_size_
;
}
inline
cl_uint
device
::
get_
max_work_item_dimensions
()
const
{
inline
cl_uint
device
::
max_work_item_dimensions
()
const
{
return
max_work_item_dimensions_
;
}
inline
size_t
device
::
get_
profiling_timer_resolution
()
const
{
inline
size_t
device
::
profiling_timer_resolution
()
const
{
return
profiling_timer_resolution_
;
}
inline
const
dim_vec
&
device
::
get_
max_work_item_sizes
()
const
{
inline
const
dim_vec
&
device
::
max_work_item_sizes
()
const
{
return
max_work_item_sizes_
;
}
inline
device_type
device
::
get_device_
type
()
const
{
inline
device_type
device
::
type
()
const
{
return
device_type_
;
}
inline
const
std
::
vector
<
std
::
string
>&
device
::
get_
extensions
()
const
{
inline
const
std
::
vector
<
std
::
string
>&
device
::
extensions
()
const
{
return
extensions_
;
}
inline
const
std
::
string
&
device
::
get_
opencl_c_version
()
const
{
inline
const
std
::
string
&
device
::
opencl_c_version
()
const
{
return
opencl_c_version_
;
}
inline
const
std
::
string
&
device
::
get_
device_vendor
()
const
{
inline
const
std
::
string
&
device
::
device_vendor
()
const
{
return
device_vendor_
;
}
inline
const
std
::
string
&
device
::
get_
device_version
()
const
{
inline
const
std
::
string
&
device
::
device_version
()
const
{
return
device_version_
;
}
inline
const
std
::
string
&
device
::
get_
driver_version
()
const
{
inline
const
std
::
string
&
device
::
driver_version
()
const
{
return
driver_version_
;
}
inline
const
std
::
string
&
device
::
get_
name
()
const
{
inline
const
std
::
string
&
device
::
name
()
const
{
return
name_
;
}
...
...
libcaf_opencl/caf/opencl/global.hpp
View file @
e5f44ee4
...
...
@@ -64,7 +64,7 @@ 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
get_
opencl_error
(
cl_int
err
);
std
::
string
opencl_error
(
cl_int
err
);
std
::
string
event_status
(
cl_event
event
);
...
...
libcaf_opencl/caf/opencl/manager.hpp
View file @
e5f44ee4
...
...
@@ -33,9 +33,10 @@
#include "caf/opencl/global.hpp"
#include "caf/opencl/program.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_actor.hpp"
#include "caf/opencl/actor_facade.hpp"
#include "caf/opencl/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
#include "caf/opencl/detail/spawn_helper.hpp"
namespace
caf
{
...
...
@@ -45,17 +46,17 @@ class manager : public actor_system::module {
public:
friend
class
program
;
friend
class
actor_system
;
friend
cl_command_queue_ptr
get_
command_queue
(
uint32_t
id
);
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
>
get
_device
(
size_t
dev_id
=
0
)
const
;
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
>
get
_device_if
(
UnaryPredicate
p
)
const
{
optional
<
device_ptr
>
find
_device_if
(
UnaryPredicate
p
)
const
{
for
(
auto
&
pl
:
platforms_
)
{
for
(
auto
&
dev
:
pl
->
get_
devices
())
{
for
(
auto
&
dev
:
pl
->
devices
())
{
if
(
p
(
dev
))
return
dev
;
}
...
...
@@ -72,7 +73,7 @@ public:
void
*
subtype_ptr
()
override
;
static
actor_system
::
module
*
make
(
actor_system
&
sys
,
caf
::
detail
::
type_list
<>
);
detail
::
type_list
<>
);
// OpenCL functionality
...
...
@@ -108,15 +109,12 @@ public:
/// @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
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
)
{
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
::
move
(
x
),
std
::
move
(
xs
)...);
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
...
...
@@ -125,16 +123,13 @@ public:
/// <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
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
)
{
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
::
move
(
x
),
std
::
move
(
xs
)...);
std
::
forward
<
T
>
(
x
),
std
::
forward
<
Ts
>
(
xs
)...);
}
/// Creates a new actor facade for an OpenCL kernel that invokes
...
...
@@ -145,7 +140,7 @@ public:
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
)
{
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
),
...
...
@@ -161,7 +156,7 @@ public:
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
)
{
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
,
...
...
@@ -176,14 +171,11 @@ public:
/// @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
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
)
{
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
)...);
...
...
@@ -195,14 +187,11 @@ public:
/// <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
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
)
{
T
&&
x
,
Ts
&&
...
xs
)
{
detail
::
cl_spawn_helper
<
false
,
Ts
...
>
f
;
return
f
(
actor_config
{
system_
.
dummy_execution_unit
()},
create_program
(
source
),
fname
,
range
,
...
...
@@ -217,14 +206,11 @@ public:
/// @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
template
<
class
Fun
,
class
...
Ts
>
typename
std
::
enable_if
<
!
opencl
::
is_opencl_arg
<
Fun
>::
value
,
actor
>::
type
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
)
{
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
),
...
...
@@ -237,14 +223,11 @@ public:
/// <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
template
<
class
Fun
,
class
...
Ts
>
typename
std
::
enable_if
<
!
opencl
::
is_opencl_arg
<
Fun
>::
value
,
actor
>::
type
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
)
{
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
,
...
...
@@ -257,14 +240,11 @@ public:
/// @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
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
)
{
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
)...);
...
...
@@ -276,14 +256,11 @@ public:
/// <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
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
)
{
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
,
...
...
libcaf_opencl/caf/opencl/mem_ref.hpp
View file @
e5f44ee4
...
...
@@ -27,13 +27,15 @@
#include "caf/optional.hpp"
#include "caf/ref_counted.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/detail/core.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace
caf
{
namespace
opencl
{
/// Updates the reference types in a message with a given event.
struct
msg_adding_event
{
msg_adding_event
(
cl
_event_ptr
event
)
:
event_
(
event
)
{
msg_adding_event
(
detail
::
raw
_event_ptr
event
)
:
event_
(
event
)
{
// nop
}
template
<
class
T
,
class
...
Ts
>
...
...
@@ -49,21 +51,24 @@ struct msg_adding_event {
ref
.
set_event
(
event_
);
return
std
::
move
(
ref
);
}
cl
_event_ptr
event_
;
detail
::
raw
_event_ptr
event_
;
};
// 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
{
class
mem_ref
:
ref_tag
{
public:
using
value_type
=
T
;
friend
struct
msg_adding_event
;
template
<
bool
PassConfig
,
class
...
Ts
>
friend
class
opencl_actor
;
friend
class
actor_facade
;
friend
class
device
;
expected
<
std
::
vector
<
T
>>
data
(
optional
<
size_t
>
result_size
=
none
)
{
...
...
@@ -85,7 +90,7 @@ public:
static_cast
<
cl_uint
>
(
prev_events
.
size
()),
prev_events
.
data
(),
&
event
);
if
(
err
!=
CL_SUCCESS
)
return
make_error
(
sec
::
runtime_error
,
get_
opencl_error
(
err
));
return
make_error
(
sec
::
runtime_error
,
opencl_error
(
err
));
// decrements the previous event we used for waiting above
event_
.
reset
(
event
,
false
);
return
buffer
;
...
...
@@ -99,11 +104,7 @@ public:
event_
.
reset
();
}
inline
const
cl_mem_ptr
&
get
()
const
{
return
memory_
;
}
inline
cl_mem_ptr
&
get
()
{
inline
const
detail
::
raw_mem_ptr
&
get
()
const
{
return
memory_
;
}
...
...
@@ -124,8 +125,9 @@ public:
// nop
}
mem_ref
(
size_t
num_elements
,
cl_command_queue_ptr
queue
,
cl_mem_ptr
memory
,
cl_mem_flags
access
,
cl_event_ptr
event
)
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
},
...
...
@@ -134,8 +136,8 @@ public:
// nop
}
mem_ref
(
size_t
num_elements
,
cl
_command_queue_ptr
queue
,
cl_mem
memory
,
cl_mem_flags
access
,
cl
_event_ptr
event
)
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
},
...
...
@@ -157,11 +159,11 @@ private:
event_
.
reset
(
e
,
increment_reference
);
}
inline
void
set_event
(
cl
_event_ptr
e
)
{
inline
void
set_event
(
detail
::
raw
_event_ptr
e
)
{
event_
=
std
::
move
(
e
);
}
inline
cl
_event_ptr
event
()
{
inline
detail
::
raw
_event_ptr
event
()
{
return
event_
;
}
...
...
@@ -171,12 +173,16 @@ private:
size_t
num_elements_
;
cl_mem_flags
access_
;
cl
_command_queue_ptr
queue_
;
cl
_event_ptr
event_
;
cl
_mem_ptr
memory_
;
detail
::
raw
_command_queue_ptr
queue_
;
detail
::
raw
_event_ptr
event_
;
detail
::
raw
_mem_ptr
memory_
;
};
}
// namespace opencl
template
<
class
T
>
struct
allowed_unsafe_message_type
<
opencl
::
mem_ref
<
T
>>
:
std
::
true_type
{};
}
// namespace caf
#endif // CAF_OPENCL_MEM_REF_HPP
libcaf_opencl/caf/opencl/platform.hpp
View file @
e5f44ee4
...
...
@@ -36,14 +36,14 @@ public:
template
<
class
T
,
class
...
Ts
>
friend
intrusive_ptr
<
T
>
caf
::
make_counted
(
Ts
&&
...);
inline
const
std
::
vector
<
device_ptr
>&
get_
devices
()
const
;
inline
const
std
::
string
&
get_
name
()
const
;
inline
const
std
::
string
&
get_
vendor
()
const
;
inline
const
std
::
string
&
get_
version
()
const
;
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
,
cl
_context_ptr
context
,
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
);
...
...
@@ -52,7 +52,7 @@ private:
static
std
::
string
platform_info
(
cl_platform_id
platform_id
,
unsigned
info_flag
);
cl_platform_id
platform_id_
;
cl
_context_ptr
context_
;
detail
::
raw
_context_ptr
context_
;
std
::
string
name_
;
std
::
string
vendor_
;
std
::
string
version_
;
...
...
@@ -63,19 +63,19 @@ private:
* implementation of inline member functions *
\******************************************************************************/
inline
const
std
::
vector
<
device_ptr
>&
platform
::
get_
devices
()
const
{
inline
const
std
::
vector
<
device_ptr
>&
platform
::
devices
()
const
{
return
devices_
;
}
inline
const
std
::
string
&
platform
::
get_
name
()
const
{
inline
const
std
::
string
&
platform
::
name
()
const
{
return
name_
;
}
inline
const
std
::
string
&
platform
::
get_
vendor
()
const
{
inline
const
std
::
string
&
platform
::
vendor
()
const
{
return
vendor_
;
}
inline
const
std
::
string
&
platform
::
get_
version
()
const
{
inline
const
std
::
string
&
platform
::
version
()
const
{
return
version_
;
}
...
...
libcaf_opencl/caf/opencl/program.hpp
View file @
e5f44ee4
...
...
@@ -27,14 +27,12 @@
#include "caf/opencl/device.hpp"
#include "caf/opencl/global.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
namespace
caf
{
namespace
opencl
{
template
<
class
...
Ts
>
class
actor_facade
;
class
program
;
using
program_ptr
=
intrusive_ptr
<
program
>
;
...
...
@@ -42,24 +40,22 @@ using program_ptr = intrusive_ptr<program>;
class
program
:
public
ref_counted
{
public:
friend
class
manager
;
template
<
class
...
Ts
>
friend
class
actor_facade
;
template
<
bool
PassConfig
,
class
...
Ts
>
friend
class
opencl_actor
;
friend
class
actor_facade
;
template
<
class
T
,
class
...
Ts
>
friend
intrusive_ptr
<
T
>
caf
::
make_counted
(
Ts
&&
...);
private:
program
(
cl_context_ptr
context
,
cl
_command_queue_ptr
queue
,
cl
_program_ptr
prog
,
std
::
map
<
std
::
string
,
cl
_kernel_ptr
>
available_kernels
);
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
();
cl
_context_ptr
context_
;
cl
_program_ptr
program_
;
cl
_command_queue_ptr
queue_
;
std
::
map
<
std
::
string
,
cl
_kernel_ptr
>
available_kernels_
;
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
...
...
libcaf_opencl/examples/scan.cpp
View file @
e5f44ee4
...
...
@@ -48,7 +48,7 @@ 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:
/// 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,
...
...
@@ -59,7 +59,6 @@ kernel void phase_1(global uint* restrict data,
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)
...
...
@@ -155,7 +154,6 @@ kernel void phase_3(global uint* restrict data,
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;
...
...
@@ -169,10 +167,7 @@ kernel void phase_3(global uint* restrict data,
}
// namespace <anonymous>
// Allow sending of unserializable references
CAF_ALLOW_UNSAFE_MESSAGE_TYPE
(
uref
);
template
<
class
T
,
class
E
=
typename
enable_if
<
is_integral
<
T
>
::
value
>::
type
>
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
;
}
...
...
@@ -182,47 +177,40 @@ int main() {
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
;
values
.
reserve
(
problem_size
);
random_device
rd
;
//Will be used to obtain a seed for the random number engine
mt19937
gen
(
rd
());
//Standard mersenne_twister_engine seeded with rd()
uvec
values
(
problem_size
);
random_device
rd
;
mt19937
gen
(
rd
());
uniform_int_distribution
<
uval
>
val_gen
(
0
,
1023
);
for
(
size_t
i
=
0
;
i
<
problem_size
;
++
i
)
values
.
emplace_back
(
val_gen
(
gen
));
std
::
generate
(
begin
(
values
),
end
(
values
),
[
&
]()
{
return
val_gen
(
gen
);
});
// ---- find device ----
auto
&
mngr
=
system
.
opencl_manager
();
//
//
string
prefix
=
"GeForce"
;
auto
opt
=
mngr
.
get
_device_if
([
&
](
const
device_ptr
dev
)
{
auto
&
name
=
dev
->
get_
name
();
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
.
get
_device
();
opt
=
mngr
.
find
_device
();
}
if
(
!
opt
)
{
cerr
<<
"Not OpenCL device available."
<<
endl
;
return
0
;
}
else
{
cerr
<<
"Found device '"
<<
(
*
opt
)
->
get_
name
()
<<
"'."
<<
endl
;
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
->
get_
max_work_group_size
()
/
2
;
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
);
};
...
...
@@ -235,7 +223,6 @@ int main() {
};
// 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
,
...
...
@@ -277,10 +264,8 @@ int main() {
in
<
uval
,
mref
>
{},
priv
<
uval
,
val
>
{}
);
// ---- composed scan actor ----
auto
scanner
=
phase3
*
phase2
*
phase1
;
// ---- scan the data ----
self
->
send
(
scanner
,
values
);
self
->
receive
(
...
...
libcaf_opencl/src/device.cpp
View file @
e5f44ee4
...
...
@@ -33,8 +33,8 @@ using namespace std;
namespace
caf
{
namespace
opencl
{
device_ptr
device
::
create
(
const
cl
_context_ptr
&
context
,
const
cl
_device_ptr
&
device_id
,
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
...
...
@@ -42,14 +42,18 @@ device_ptr device::create(const cl_context_ptr& context,
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
;
properties
|=
out_of_order
?
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
:
0
;
if
(
out_of_order
)
properties
|=
CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE
;
// create the command queue
cl_command_queue_ptr
command_queue
{
v2get
(
CAF_CLF
(
clCreateCommandQueue
),
context
.
get
(),
device_id
.
get
(),
properties
),
false
};
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
);
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
);
...
...
@@ -103,7 +107,8 @@ void device::synchronize() {
clFinish
(
queue_
.
get
());
}
string
device
::
info_string
(
const
cl_device_ptr
&
device_id
,
unsigned
info_flag
)
{
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
);
...
...
@@ -112,8 +117,10 @@ string device::info_string(const cl_device_ptr& device_id, unsigned info_flag) {
return
string
(
buffer
.
data
());
}
device
::
device
(
cl_device_ptr
device_id
,
cl_command_queue_ptr
queue
,
cl_context_ptr
context
,
unsigned
id
)
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
)),
...
...
libcaf_opencl/src/global.cpp
View file @
e5f44ee4
...
...
@@ -57,7 +57,7 @@ device_type device_type_from_ulong(cl_ulong dev) {
}
}
std
::
string
get_
opencl_error
(
cl_int
err
)
{
std
::
string
opencl_error
(
cl_int
err
)
{
switch
(
err
)
{
case
CL_SUCCESS
:
return
"CL_SUCCESS"
;
...
...
libcaf_opencl/src/manager.cpp
View file @
e5f44ee4
...
...
@@ -25,23 +25,24 @@
#include "caf/opencl/device.hpp"
#include "caf/opencl/manager.hpp"
#include "caf/opencl/platform.hpp"
#include "caf/opencl/smart_ptr.hpp"
#include "caf/opencl/opencl_err.hpp"
#include "caf/opencl/detail/raw_ptr.hpp"
using
namespace
std
;
namespace
caf
{
namespace
opencl
{
optional
<
device_ptr
>
manager
::
get
_device
(
size_t
dev_id
)
const
{
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
->
get_
devices
().
size
();
to
+=
pl
->
devices
().
size
();
if
(
dev_id
>=
from
&&
dev_id
<
to
)
return
pl
->
get_
devices
()[
dev_id
-
from
];
return
pl
->
devices
()[
dev_id
-
from
];
}
return
none
;
}
...
...
@@ -59,7 +60,7 @@ void manager::init(actor_system_config&) {
for
(
auto
&
pl_id
:
platform_ids
)
{
platforms_
.
push_back
(
platform
::
create
(
pl_id
,
current_device_id
));
current_device_id
+=
static_cast
<
unsigned
>
(
platforms_
.
back
()
->
get_
devices
().
size
());
static_cast
<
unsigned
>
(
platforms_
.
back
()
->
devices
().
size
());
}
}
...
...
@@ -108,7 +109,7 @@ program_ptr manager::create_program_from_file(const char* path,
program_ptr
manager
::
create_program
(
const
char
*
kernel_source
,
const
char
*
options
,
uint32_t
device_id
)
{
auto
dev
=
get
_device
(
device_id
);
auto
dev
=
find
_device
(
device_id
);
if
(
!
dev
)
{
ostringstream
oss
;
oss
<<
"No device with id '"
<<
device_id
<<
"' found."
;
...
...
@@ -144,7 +145,7 @@ program_ptr manager::create_program(const char* kernel_source,
const
device_ptr
dev
)
{
// create program object from kernel source
size_t
kernel_source_length
=
strlen
(
kernel_source
);
cl
_program_ptr
pptr
;
detail
::
raw
_program_ptr
pptr
;
pptr
.
reset
(
v2get
(
CAF_CLF
(
clCreateProgramWithSource
),
dev
->
context_
.
get
(),
1u
,
&
kernel_source
,
&
kernel_source_length
),
false
);
...
...
@@ -153,7 +154,7 @@ program_ptr manager::create_program(const char* kernel_source,
auto
err
=
clBuildProgram
(
pptr
.
get
(),
1
,
&
dev_tmp
,
options
,
nullptr
,
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
oss
<<
"clBuildProgram: "
<<
get_
opencl_error
(
err
);
oss
<<
"clBuildProgram: "
<<
opencl_error
(
err
);
if
(
err
==
CL_BUILD_PROGRAM_FAILURE
)
{
size_t
buildlog_buffer_size
=
0
;
// get the log length
...
...
@@ -180,14 +181,14 @@ program_ptr manager::create_program(const char* kernel_source,
}
cl_uint
number_of_kernels
=
0
;
clCreateKernelsInProgram
(
pptr
.
get
(),
0u
,
nullptr
,
&
number_of_kernels
);
map
<
string
,
cl
_kernel_ptr
>
available_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
)
{
ostringstream
oss
;
oss
<<
"clCreateKernelsInProgram: "
<<
get_
opencl_error
(
err
);
oss
<<
"clCreateKernelsInProgram: "
<<
opencl_error
(
err
);
throw
runtime_error
(
oss
.
str
());
}
for
(
cl_uint
i
=
0
;
i
<
number_of_kernels
;
++
i
)
{
...
...
@@ -199,10 +200,10 @@ program_ptr manager::create_program(const char* kernel_source,
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
oss
<<
"clGetKernelInfo (CL_KERNEL_FUNCTION_NAME): "
<<
get_
opencl_error
(
err
);
<<
opencl_error
(
err
);
throw
runtime_error
(
oss
.
str
());
}
cl
_kernel_ptr
kernel
;
detail
::
raw
_kernel_ptr
kernel
;
kernel
.
reset
(
move
(
kernels
[
i
]));
available_kernels
.
emplace
(
string
(
name
.
data
()),
move
(
kernel
));
}
...
...
libcaf_opencl/src/opencl_err.cpp
View file @
e5f44ee4
...
...
@@ -27,7 +27,7 @@ void throwcl(const char* fname, cl_int err) {
if
(
err
!=
CL_SUCCESS
)
{
std
::
string
errstr
=
fname
;
errstr
+=
": "
;
errstr
+=
get_
opencl_error
(
err
);
errstr
+=
opencl_error
(
err
);
throw
std
::
runtime_error
(
std
::
move
(
errstr
));
}
}
...
...
libcaf_opencl/src/platform.cpp
View file @
e5f44ee4
...
...
@@ -48,11 +48,13 @@ platform_ptr platform::create(cl_platform_id platform_id,
v2callcl
(
CAF_CLF
(
clGetDeviceIDs
),
platform_id
,
device_type
,
discoverd
,
(
ids
.
data
()
+
known
));
}
vector
<
cl
_device_ptr
>
devices
;
vector
<
detail
::
raw
_device_ptr
>
devices
;
devices
.
resize
(
ids
.
size
());
auto
lift
=
[](
cl_device_id
ptr
)
{
return
cl_device_ptr
{
ptr
,
false
};
};
auto
lift
=
[](
cl_device_id
ptr
)
{
return
detail
::
raw_device_ptr
{
ptr
,
false
};
};
transform
(
ids
.
begin
(),
ids
.
end
(),
devices
.
begin
(),
lift
);
cl
_context_ptr
context
;
detail
::
raw
_context_ptr
context
;
context
.
reset
(
v2get
(
CAF_CLF
(
clCreateContext
),
nullptr
,
static_cast
<
unsigned
>
(
ids
.
size
()),
ids
.
data
(),
pfn_notify
,
nullptr
),
...
...
@@ -87,7 +89,7 @@ string platform::platform_info(cl_platform_id platform_id,
return
string
(
buffer
.
data
());
}
platform
::
platform
(
cl_platform_id
platform_id
,
cl
_context_ptr
context
,
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
),
...
...
libcaf_opencl/src/program.cpp
View file @
e5f44ee4
...
...
@@ -32,9 +32,10 @@ using namespace std;
namespace
caf
{
namespace
opencl
{
program
::
program
(
cl_context_ptr
context
,
cl_command_queue_ptr
queue
,
cl_program_ptr
prog
,
map
<
string
,
cl_kernel_ptr
>
available_kernels
)
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
)),
...
...
libcaf_opencl/test/opencl.cpp
View file @
e5f44ee4
...
...
@@ -16,11 +16,11 @@ 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
;
// required to allow sending mem_ref<int> in messages
CAF_ALLOW_UNSAFE_MESSAGE_TYPE
(
mem_ref
<
int
>
);
namespace
{
using
ivec
=
vector
<
int
>
;
...
...
@@ -327,7 +327,7 @@ void check_mref_results(const string& description,
void
test_opencl
(
actor_system
&
sys
)
{
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device_if
([](
const
device_ptr
){
return
true
;
});
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
);
...
...
@@ -425,7 +425,7 @@ void test_opencl(actor_system& sys) {
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
get_
max_work_item_sizes
()[
0
],
size_t
{
512
});
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
;
...
...
@@ -475,7 +475,7 @@ void test_opencl(actor_system& sys) {
void
test_arguments
(
actor_system
&
sys
)
{
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device_if
([](
const
device_ptr
){
return
true
;
});
auto
opt
=
mngr
.
find
_device_if
([](
const
device_ptr
){
return
true
;
});
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
scoped_actor
self
{
sys
};
...
...
@@ -598,7 +598,7 @@ CAF_TEST(opencl_mem_refs) {
cfg
.
load
<
opencl
::
manager
>
();
actor_system
system
{
cfg
};
auto
&
mngr
=
system
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
// global arguments
...
...
@@ -630,44 +630,44 @@ CAF_TEST(opencl_mem_refs) {
CAF_TEST
(
opencl_argument_info
)
{
using
base_t
=
int
;
using
in_arg_t
=
caf
::
detail
::
type_list
<
opencl
::
in
<
base_t
>>
;
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
caf
::
detail
::
tl_head
<
in_arg_info_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
=
caf
::
detail
::
type_list
<
opencl
::
out
<
base_t
>>
;
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
caf
::
detail
::
tl_head
<
out_arg_info_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
=
caf
::
detail
::
type_list
<
opencl
::
in_out
<
base_t
>>
;
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
caf
::
detail
::
tl_head
<
io_arg_info_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
=
caf
::
detail
::
type_list
<
opencl
::
in
<
base_t
>
,
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
caf
::
detail
::
tl_at
<
arg_info_list_t
,
0
>::
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
caf
::
detail
::
tl_at
<
arg_info_list_t
,
1
>::
type
;
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
caf
::
detail
::
tl_at
<
arg_info_list_t
,
2
>::
type
;
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
caf
::
detail
::
tl_at
<
arg_info_list_t
,
3
>::
type
;
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
caf
::
detail
::
tl_at
<
arg_info_list_t
,
4
>::
type
;
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
caf
::
detail
::
tl_at
<
arg_info_list_t
,
5
>::
type
;
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.
...
...
@@ -677,7 +677,7 @@ CAF_TEST(opencl_argument_info) {
void
test_in_val_out_val
(
actor_system
&
sys
)
{
CAF_MESSAGE
(
"Testing in: val -> out: val "
);
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -733,7 +733,7 @@ void test_in_val_out_val(actor_system& sys) {
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
get_
max_work_item_sizes
()[
0
],
size_t
{
512
});
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
;
...
...
@@ -775,7 +775,7 @@ 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
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -803,7 +803,7 @@ void test_in_val_out_mref(actor_system& sys) {
" (kernel passed directly)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
get_
max_work_item_sizes
()[
0
],
size_t
{
512
});
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
;
...
...
@@ -830,7 +830,7 @@ 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
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -860,7 +860,7 @@ void test_in_mref_out_val(actor_system& sys) {
" (kernel passed directly)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
get_
max_work_item_sizes
()[
0
],
size_t
{
512
});
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
;
...
...
@@ -888,7 +888,7 @@ 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
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -919,7 +919,7 @@ void test_in_mref_out_mref(actor_system& sys) {
" (kernel passed directly)"
,
res1
,
result
);
},
others
>>
wrong_msg
);
// test for manuel return size selection (max workgroup size 1d)
auto
max_wg_size
=
min
(
dev
->
get_
max_work_item_sizes
()[
0
],
size_t
{
512
});
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
;
...
...
@@ -948,7 +948,7 @@ void test_varying_arguments(actor_system& sys) {
"(Might fail on some integrated GPUs)"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -983,7 +983,7 @@ void test_inout(actor_system& sys) {
CAF_MESSAGE
(
"Testing in_out arguments"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -1029,7 +1029,7 @@ void test_priv(actor_system& sys) {
CAF_MESSAGE
(
"Testing priv argument"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -1062,7 +1062,7 @@ void test_local(actor_system& sys) {
CAF_MESSAGE
(
"Testing local argument"
);
// setup
auto
&
mngr
=
sys
.
opencl_manager
();
auto
opt
=
mngr
.
get
_device
(
0
);
auto
opt
=
mngr
.
find
_device
(
0
);
CAF_REQUIRE
(
opt
);
auto
dev
=
*
opt
;
auto
prog
=
mngr
.
create_program
(
kernel_source
,
""
,
dev
);
...
...
@@ -1098,7 +1098,7 @@ void test_local(actor_system& sys) {
},
others
>>
wrong_msg
);
}
CAF_TEST
(
opencl_opencl_actor
)
{
CAF_TEST
(
actor_facade
)
{
actor_system_config
cfg
;
cfg
.
load
<
opencl
::
manager
>
()
.
add_message_type
<
ivec
>
(
"int_vector"
)
...
...
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