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
f0a23d1b
Commit
f0a23d1b
authored
Jun 11, 2015
by
Dominik Charousset
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
Reformatted according to new coding style
parent
c2c73d90
Changes
11
Hide whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
158 additions
and
174 deletions
+158
-174
libcaf_opencl/caf/opencl/actor_facade.hpp
libcaf_opencl/caf/opencl/actor_facade.hpp
+35
-35
libcaf_opencl/caf/opencl/command.hpp
libcaf_opencl/caf/opencl/command.hpp
+35
-35
libcaf_opencl/caf/opencl/device_info.hpp
libcaf_opencl/caf/opencl/device_info.hpp
+15
-15
libcaf_opencl/caf/opencl/global.hpp
libcaf_opencl/caf/opencl/global.hpp
+1
-3
libcaf_opencl/caf/opencl/opencl_metainfo.hpp
libcaf_opencl/caf/opencl/opencl_metainfo.hpp
+5
-7
libcaf_opencl/caf/opencl/program.hpp
libcaf_opencl/caf/opencl/program.hpp
+9
-13
libcaf_opencl/caf/opencl/spawn_cl.hpp
libcaf_opencl/caf/opencl/spawn_cl.hpp
+18
-26
libcaf_opencl/examples/proper_matrix.cpp
libcaf_opencl/examples/proper_matrix.cpp
+14
-14
libcaf_opencl/src/opencl_metainfo.cpp
libcaf_opencl/src/opencl_metainfo.cpp
+5
-5
libcaf_opencl/src/program.cpp
libcaf_opencl/src/program.cpp
+6
-6
libcaf_opencl/test/opencl.cpp
libcaf_opencl/test/opencl.cpp
+15
-15
No files found.
libcaf_opencl/caf/opencl/actor_facade.hpp
View file @
f0a23d1b
...
@@ -54,7 +54,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -54,7 +54,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
friend
class
command
<
actor_facade
,
Ret
>
;
friend
class
command
<
actor_facade
,
Ret
>
;
public:
public:
using
arg_types
=
detail
::
type_list
<
typename
std
::
decay
<
Args
>::
type
...
>
;
using
arg_types
=
detail
::
type_list
<
typename
std
::
decay
<
Args
>::
type
...
>
;
using
arg_mapping
=
std
::
function
<
optional
<
message
>
(
message
&
)
>
;
using
arg_mapping
=
std
::
function
<
optional
<
message
>
(
message
&
)
>
;
using
result_mapping
=
std
::
function
<
message
(
Ret
&
)
>
;
using
result_mapping
=
std
::
function
<
message
(
Ret
&
)
>
;
...
@@ -74,7 +74,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -74,7 +74,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
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
()
!=
global_dims
.
size
())
{
if
(
!
vec
.
empty
()
&&
vec
.
size
()
!=
global_dims
.
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"
;
...
@@ -85,7 +85,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -85,7 +85,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
check_vec
(
offsets
,
"offsets"
);
check_vec
(
offsets
,
"offsets"
);
check_vec
(
local_dims
,
"local dimensions"
);
check_vec
(
local_dims
,
"local dimensions"
);
kernel_ptr
kernel
;
kernel_ptr
kernel
;
kernel
.
reset
(
v2get
(
CAF_CLF
(
clCreateKernel
),
prog
.
m_program
.
get
(),
kernel
.
reset
(
v2get
(
CAF_CLF
(
clCreateKernel
),
prog
.
program_
.
get
(),
kernel_name
),
kernel_name
),
false
);
false
);
if
(
result_size
==
0
)
{
if
(
result_size
==
0
)
{
...
@@ -100,45 +100,45 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -100,45 +100,45 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
void
enqueue
(
const
actor_addr
&
sender
,
message_id
mid
,
message
content
,
void
enqueue
(
const
actor_addr
&
sender
,
message_id
mid
,
message
content
,
execution_unit
*
)
override
{
execution_unit
*
)
override
{
CAF_LOG_TRACE
(
""
);
CAF_LOG_TRACE
(
""
);
if
(
m
_map_args
)
{
if
(
m
ap_args_
)
{
auto
mapped
=
m
_map_args
(
content
);
auto
mapped
=
m
ap_args_
(
content
);
if
(
!
mapped
)
{
if
(
!
mapped
)
{
return
;
return
;
}
}
content
=
std
::
move
(
*
mapped
);
content
=
std
::
move
(
*
mapped
);
}
}
typename
detail
::
il_indices
<
arg_types
>::
type
indices
;
typename
detail
::
il_indices
<
arg_types
>::
type
indices
;
if
(
!
content
.
match_elements
(
arg_types
{}))
{
if
(
!
content
.
match_elements
(
arg_types
{}))
{
return
;
return
;
}
}
response_promise
handle
{
this
->
address
(),
sender
,
mid
.
response_id
()};
response_promise
handle
{
this
->
address
(),
sender
,
mid
.
response_id
()};
evnt_vec
events
;
evnt_vec
events
;
args_vec
arguments
;
args_vec
arguments
;
add_arguments_to_kernel
<
Ret
>
(
events
,
arguments
,
m_result_size
,
add_arguments_to_kernel
<
Ret
>
(
events
,
arguments
,
result_size_
,
content
,
indices
);
content
,
indices
);
auto
cmd
=
make_counted
<
command_type
>
(
handle
,
this
,
auto
cmd
=
make_counted
<
command_type
>
(
handle
,
this
,
std
::
move
(
events
),
std
::
move
(
events
),
std
::
move
(
arguments
),
std
::
move
(
arguments
),
m_result_size
,
result_size_
,
std
::
move
(
content
));
std
::
move
(
content
));
cmd
->
enqueue
();
cmd
->
enqueue
();
}
}
private:
private:
actor_facade
(
const
program
&
prog
,
kernel_ptr
kernel
,
actor_facade
(
const
program
&
prog
,
kernel_ptr
kernel
,
const
dim_vec
&
global_dimensions
,
const
dim_vec
&
global_offsets
,
const
dim_vec
&
global_dimensions
,
const
dim_vec
&
global_offsets
,
const
dim_vec
&
local_dimensions
,
size_t
result_size
,
const
dim_vec
&
local_dimensions
,
size_t
result_size
,
arg_mapping
map_args
,
result_mapping
map_result
)
arg_mapping
map_args
,
result_mapping
map_result
)
:
m_kernel
(
kernel
),
:
kernel_
(
kernel
),
m_program
(
prog
.
m_program
),
program_
(
prog
.
program_
),
m_context
(
prog
.
m_context
),
context_
(
prog
.
context_
),
m_queue
(
prog
.
m_queue
),
queue_
(
prog
.
queue_
),
m_global_dimensions
(
global_dimensions
),
global_dimensions_
(
global_dimensions
),
m_global_offsets
(
global_offsets
),
global_offsets_
(
global_offsets
),
m_local_dimensions
(
local_dimensions
),
local_dimensions_
(
local_dimensions
),
m_result_size
(
result_size
),
result_size_
(
result_size
),
m
_map_args
(
std
::
move
(
map_args
)),
m
ap_args_
(
std
::
move
(
map_args
)),
m
_map_result
(
std
::
move
(
map_result
))
{
m
ap_result_
(
std
::
move
(
map_result
))
{
CAF_LOG_TRACE
(
"id: "
<<
this
->
id
());
CAF_LOG_TRACE
(
"id: "
<<
this
->
id
());
}
}
...
@@ -147,10 +147,10 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -147,10 +147,10 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
// rotate left (output buffer to the end)
// rotate left (output buffer to the end)
std
::
rotate
(
arguments
.
begin
(),
arguments
.
begin
()
+
1
,
arguments
.
end
());
std
::
rotate
(
arguments
.
begin
(),
arguments
.
begin
()
+
1
,
arguments
.
end
());
for
(
cl_uint
i
=
0
;
i
<
arguments
.
size
();
++
i
)
{
for
(
cl_uint
i
=
0
;
i
<
arguments
.
size
();
++
i
)
{
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
m_kernel
.
get
(),
i
,
v1callcl
(
CAF_CLF
(
clSetKernelArg
),
kernel_
.
get
(),
i
,
sizeof
(
cl_mem
),
static_cast
<
void
*>
(
&
arguments
[
i
]));
sizeof
(
cl_mem
),
static_cast
<
void
*>
(
&
arguments
[
i
]));
}
}
clFlush
(
m_queue
.
get
());
clFlush
(
queue_
.
get
());
}
}
template
<
long
I
,
long
...
Is
>
template
<
long
I
,
long
...
Is
>
...
@@ -159,10 +159,10 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -159,10 +159,10 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
using
value_type
=
typename
detail
::
tl_at
<
arg_types
,
I
>::
type
;
using
value_type
=
typename
detail
::
tl_at
<
arg_types
,
I
>::
type
;
auto
&
arg
=
msg
.
get_as
<
value_type
>
(
I
);
auto
&
arg
=
msg
.
get_as
<
value_type
>
(
I
);
size_t
buffer_size
=
sizeof
(
value_type
)
*
arg
.
size
();
size_t
buffer_size
=
sizeof
(
value_type
)
*
arg
.
size
();
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
m_context
.
get
(),
auto
buffer
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
cl_mem_flags
{
CL_MEM_READ_ONLY
},
buffer_size
,
nullptr
);
cl_mem_flags
{
CL_MEM_READ_ONLY
},
buffer_size
,
nullptr
);
cl_event
event
=
v1get
<
cl_event
>
(
CAF_CLF
(
clEnqueueWriteBuffer
),
cl_event
event
=
v1get
<
cl_event
>
(
CAF_CLF
(
clEnqueueWriteBuffer
),
m_queue
.
get
(),
buffer
,
cl_bool
{
CL_FALSE
},
queue_
.
get
(),
buffer
,
cl_bool
{
CL_FALSE
},
cl_uint
{
0
},
buffer_size
,
arg
.
data
());
cl_uint
{
0
},
buffer_size
,
arg
.
data
());
events
.
push_back
(
std
::
move
(
event
));
events
.
push_back
(
std
::
move
(
event
));
mem_ptr
tmp
;
mem_ptr
tmp
;
...
@@ -176,7 +176,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -176,7 +176,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
void
add_arguments_to_kernel
(
evnt_vec
&
events
,
args_vec
&
arguments
,
void
add_arguments_to_kernel
(
evnt_vec
&
events
,
args_vec
&
arguments
,
size_t
ret_size
,
message
&
msg
,
Token
tk
)
{
size_t
ret_size
,
message
&
msg
,
Token
tk
)
{
arguments
.
clear
();
arguments
.
clear
();
auto
buf
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
m_context
.
get
(),
auto
buf
=
v2get
(
CAF_CLF
(
clCreateBuffer
),
context_
.
get
(),
cl_mem_flags
{
CL_MEM_WRITE_ONLY
},
cl_mem_flags
{
CL_MEM_WRITE_ONLY
},
sizeof
(
typename
R
::
value_type
)
*
ret_size
,
nullptr
);
sizeof
(
typename
R
::
value_type
)
*
ret_size
,
nullptr
);
mem_ptr
tmp
;
mem_ptr
tmp
;
...
@@ -185,16 +185,16 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -185,16 +185,16 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
add_arguments_to_kernel_rec
(
events
,
arguments
,
msg
,
tk
);
add_arguments_to_kernel_rec
(
events
,
arguments
,
msg
,
tk
);
}
}
kernel_ptr
m_kernel
;
kernel_ptr
kernel_
;
program_ptr
m_program
;
program_ptr
program_
;
context_ptr
m_context
;
context_ptr
context_
;
command_queue_ptr
m_queue
;
command_queue_ptr
queue_
;
dim_vec
m_global_dimensions
;
dim_vec
global_dimensions_
;
dim_vec
m_global_offsets
;
dim_vec
global_offsets_
;
dim_vec
m_local_dimensions
;
dim_vec
local_dimensions_
;
size_t
m_result_size
;
size_t
result_size_
;
arg_mapping
m
_map_args
;
arg_mapping
m
ap_args_
;
result_mapping
m
_map_result
;
result_mapping
m
ap_result_
;
};
};
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/caf/opencl/command.hpp
View file @
f0a23d1b
...
@@ -38,23 +38,23 @@ namespace opencl {
...
@@ -38,23 +38,23 @@ namespace opencl {
template
<
typename
T
,
typename
R
>
template
<
typename
T
,
typename
R
>
class
command
:
public
ref_counted
{
class
command
:
public
ref_counted
{
public:
public:
command
(
response_promise
handle
,
intrusive_ptr
<
T
>
actor_facade
,
command
(
response_promise
handle
,
intrusive_ptr
<
T
>
actor_facade
,
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
mem_ptr
>
arguments
,
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
mem_ptr
>
arguments
,
size_t
result_size
,
message
msg
)
size_t
result_size
,
message
msg
)
:
m_result_size
(
result_size
),
:
result_size_
(
result_size
),
m_handle
(
handle
),
handle_
(
handle
),
m_actor_facade
(
actor_facade
),
actor_facade_
(
actor_facade
),
m_queue
(
actor_facade
->
m_queue
),
queue_
(
actor_facade
->
queue_
),
m_events
(
std
::
move
(
events
)),
events_
(
std
::
move
(
events
)),
m_arguments
(
std
::
move
(
arguments
)),
arguments_
(
std
::
move
(
arguments
)),
m_result
(
result_size
),
result_
(
result_size
),
m
_msg
(
msg
)
{
m
sg_
(
msg
)
{
// nop
// nop
}
}
~
command
()
{
~
command
()
{
for
(
auto
&
e
:
m_events
)
{
for
(
auto
&
e
:
events_
)
{
v1callcl
(
CAF_CLF
(
clReleaseEvent
),
e
);
v1callcl
(
CAF_CLF
(
clReleaseEvent
),
e
);
}
}
}
}
...
@@ -70,13 +70,13 @@ class command : public ref_counted {
...
@@ -70,13 +70,13 @@ class command : public ref_counted {
};
};
// 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
(
m_queue
.
get
(),
m_actor_facade
->
m_kernel
.
get
(),
queue_
.
get
(),
actor_facade_
->
kernel_
.
get
(),
static_cast
<
cl_uint
>
(
m_actor_facade
->
m_global_dimensions
.
size
()),
static_cast
<
cl_uint
>
(
actor_facade_
->
global_dimensions_
.
size
()),
data_or_nullptr
(
m_actor_facade
->
m_global_offsets
),
data_or_nullptr
(
actor_facade_
->
global_offsets_
),
data_or_nullptr
(
m_actor_facade
->
m_global_dimensions
),
data_or_nullptr
(
actor_facade_
->
global_dimensions_
),
data_or_nullptr
(
m_actor_facade
->
m_local_dimensions
),
data_or_nullptr
(
actor_facade_
->
local_dimensions_
),
static_cast
<
cl_uint
>
(
m_events
.
size
()),
static_cast
<
cl_uint
>
(
events_
.
size
()),
(
m_events
.
empty
()
?
nullptr
:
m_events
.
data
()),
&
event_k
);
(
events_
.
empty
()
?
nullptr
:
events_
.
data
()),
&
event_k
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOGMF
(
CAF_ERROR
,
"clEnqueueNDRangeKernel: "
<<
get_opencl_error
(
err
));
CAF_LOGMF
(
CAF_ERROR
,
"clEnqueueNDRangeKernel: "
<<
get_opencl_error
(
err
));
this
->
deref
();
// or can anything actually happen?
this
->
deref
();
// or can anything actually happen?
...
@@ -84,9 +84,9 @@ class command : public ref_counted {
...
@@ -84,9 +84,9 @@ class command : public ref_counted {
}
else
{
}
else
{
cl_event
event_r
;
cl_event
event_r
;
err
=
err
=
clEnqueueReadBuffer
(
m_queue
.
get
(),
m_arguments
.
back
().
get
(),
CL_FALSE
,
clEnqueueReadBuffer
(
queue_
.
get
(),
arguments_
.
back
().
get
(),
CL_FALSE
,
0
,
sizeof
(
typename
R
::
value_type
)
*
m_result_size
,
0
,
sizeof
(
typename
R
::
value_type
)
*
result_size_
,
m_result
.
data
(),
1
,
&
event_k
,
&
event_r
);
result_
.
data
(),
1
,
&
event_k
,
&
event_r
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
this
->
deref
();
// failed to enqueue command
this
->
deref
();
// failed to enqueue command
throw
std
::
runtime_error
(
"clEnqueueReadBuffer: "
+
throw
std
::
runtime_error
(
"clEnqueueReadBuffer: "
+
...
@@ -105,29 +105,29 @@ class command : public ref_counted {
...
@@ -105,29 +105,29 @@ class command : public ref_counted {
return
;
return
;
}
}
err
=
clFlush
(
m_queue
.
get
());
err
=
clFlush
(
queue_
.
get
());
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CAF_LOGMF
(
CAF_ERROR
,
"clFlush: "
<<
get_opencl_error
(
err
));
CAF_LOGMF
(
CAF_ERROR
,
"clFlush: "
<<
get_opencl_error
(
err
));
}
}
m_events
.
push_back
(
std
::
move
(
event_k
));
events_
.
push_back
(
std
::
move
(
event_k
));
m_events
.
push_back
(
std
::
move
(
event_r
));
events_
.
push_back
(
std
::
move
(
event_r
));
}
}
}
}
private:
private:
size_t
m_result_size
;
size_t
result_size_
;
response_promise
m_handle
;
response_promise
handle_
;
intrusive_ptr
<
T
>
m_actor_facade
;
intrusive_ptr
<
T
>
actor_facade_
;
command_queue_ptr
m_queue
;
command_queue_ptr
queue_
;
std
::
vector
<
cl_event
>
m_events
;
std
::
vector
<
cl_event
>
events_
;
std
::
vector
<
mem_ptr
>
m_arguments
;
std
::
vector
<
mem_ptr
>
arguments_
;
R
m_result
;
R
result_
;
message
m
_msg
;
// required to keep the argument buffers alive (async copy)
message
m
sg_
;
// required to keep the argument buffers alive (async copy)
void
handle_results
()
{
void
handle_results
()
{
auto
&
map_fun
=
m_actor_facade
->
m_map_result
;
auto
&
map_fun
=
actor_facade_
->
map_result_
;
auto
msg
=
map_fun
?
map_fun
(
m_result
)
:
make_message
(
std
::
move
(
m_result
));
auto
msg
=
map_fun
?
map_fun
(
result_
)
:
make_message
(
std
::
move
(
result_
));
m_handle
.
deliver
(
std
::
move
(
msg
));
handle_
.
deliver
(
std
::
move
(
msg
));
}
}
};
};
...
...
libcaf_opencl/caf/opencl/device_info.hpp
View file @
f0a23d1b
...
@@ -31,26 +31,26 @@ class device_info {
...
@@ -31,26 +31,26 @@ class device_info {
friend
class
program
;
friend
class
program
;
public:
public:
device_info
(
device_ptr
device
,
command_queue_ptr
queue
,
device_info
(
device_ptr
device
,
command_queue_ptr
queue
,
size_t
work_group_size
,
cl_uint
dimensons
,
size_t
work_group_size
,
cl_uint
dimensons
,
const
dim_vec
&
items_per_dimension
)
const
dim_vec
&
items_per_dimension
)
:
m
_max_work_group_size
(
work_group_size
),
:
m
ax_work_group_size_
(
work_group_size
),
m
_max_dimensions
(
dimensons
),
m
ax_dimensions_
(
dimensons
),
m
_max_work_items_per_dim
(
items_per_dimension
),
m
ax_work_items_per_dim_
(
items_per_dimension
),
m_device
(
device
),
device_
(
device
),
m_cmd_queue
(
queue
)
{}
cmd_queue_
(
queue
)
{}
inline
size_t
get_max_work_group_size
();
inline
size_t
get_max_work_group_size
();
inline
cl_uint
get_max_dimensions
();
inline
cl_uint
get_max_dimensions
();
inline
dim_vec
get_max_work_items_per_dim
();
inline
dim_vec
get_max_work_items_per_dim
();
private:
private:
size_t
m
_max_work_group_size
;
size_t
m
ax_work_group_size_
;
cl_uint
m
_max_dimensions
;
cl_uint
m
ax_dimensions_
;
dim_vec
m
_max_work_items_per_dim
;
dim_vec
m
ax_work_items_per_dim_
;
device_ptr
m_device
;
device_ptr
device_
;
command_queue_ptr
m_cmd_queue
;
command_queue_ptr
cmd_queue_
;
};
};
/******************************************************************************\
/******************************************************************************\
...
@@ -58,13 +58,13 @@ class device_info {
...
@@ -58,13 +58,13 @@ class device_info {
\******************************************************************************/
\******************************************************************************/
inline
size_t
device_info
::
get_max_work_group_size
()
{
inline
size_t
device_info
::
get_max_work_group_size
()
{
return
m
_max_work_group_size
;
return
m
ax_work_group_size_
;
}
}
inline
cl_uint
device_info
::
get_max_dimensions
()
{
return
m
_max_dimensions
;
}
inline
cl_uint
device_info
::
get_max_dimensions
()
{
return
m
ax_dimensions_
;
}
inline
dim_vec
device_info
::
get_max_work_items_per_dim
()
{
inline
dim_vec
device_info
::
get_max_work_items_per_dim
()
{
return
m
_max_work_items_per_dim
;
return
m
ax_work_items_per_dim_
;
}
}
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/caf/opencl/global.hpp
View file @
f0a23d1b
...
@@ -40,9 +40,7 @@ cl_int clRetainDeviceDummy(cl_device_id);
...
@@ -40,9 +40,7 @@ cl_int clRetainDeviceDummy(cl_device_id);
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
/**
/// A vector of up to three elements used for OpenCL dimensions.
* A vector of up to three elements used for OpenCL dimensions.
*/
using
dim_vec
=
detail
::
limited_vector
<
size_t
,
3
>
;
using
dim_vec
=
detail
::
limited_vector
<
size_t
,
3
>
;
std
::
string
get_opencl_error
(
cl_int
err
);
std
::
string
get_opencl_error
(
cl_int
err
);
...
...
libcaf_opencl/caf/opencl/opencl_metainfo.hpp
View file @
f0a23d1b
...
@@ -46,23 +46,21 @@ class opencl_metainfo : public detail::abstract_singleton {
...
@@ -46,23 +46,21 @@ class opencl_metainfo : public detail::abstract_singleton {
friend
class
detail
::
singletons
;
friend
class
detail
::
singletons
;
friend
command_queue_ptr
get_command_queue
(
uint32_t
id
);
friend
command_queue_ptr
get_command_queue
(
uint32_t
id
);
public:
public:
const
std
::
vector
<
device_info
>
get_devices
()
const
;
const
std
::
vector
<
device_info
>
get_devices
()
const
;
/**
/// Get opencl_metainfo instance.
* Get opencl_metainfo instance.
*/
static
opencl_metainfo
*
instance
();
static
opencl_metainfo
*
instance
();
private:
private:
opencl_metainfo
()
=
default
;
opencl_metainfo
()
=
default
;
void
stop
()
override
;
void
stop
()
override
;
void
initialize
()
override
;
void
initialize
()
override
;
void
dispose
()
override
;
void
dispose
()
override
;
context_ptr
m_context
;
context_ptr
context_
;
std
::
vector
<
device_info
>
m_devices
;
std
::
vector
<
device_info
>
devices_
;
};
};
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/caf/opencl/program.hpp
View file @
f0a23d1b
...
@@ -31,29 +31,25 @@ namespace opencl {
...
@@ -31,29 +31,25 @@ namespace opencl {
template
<
typename
Signature
>
template
<
typename
Signature
>
class
actor_facade
;
class
actor_facade
;
/**
/// @brief A wrapper for OpenCL's cl_program.
* @brief A wrapper for OpenCL's cl_program.
*/
class
program
{
class
program
{
template
<
typename
Signature
>
template
<
typename
Signature
>
friend
class
actor_facade
;
friend
class
actor_facade
;
public:
public:
/**
/// @brief Factory method, that creates a caf::opencl::program
* @brief Factory method, that creates a caf::opencl::program
/// from a given @p kernel_source.
* from a given @p kernel_source.
/// @returns A program object.
* @returns A program object.
*/
static
program
create
(
const
char
*
kernel_source
,
static
program
create
(
const
char
*
kernel_source
,
const
char
*
options
=
nullptr
,
uint32_t
device_id
=
0
);
const
char
*
options
=
nullptr
,
uint32_t
device_id
=
0
);
private:
private:
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program_ptr
program
);
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program_ptr
program
);
context_ptr
m_context
;
context_ptr
context_
;
program_ptr
m_program
;
program_ptr
program_
;
command_queue_ptr
m_queue
;
command_queue_ptr
queue_
;
};
};
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/caf/opencl/spawn_cl.hpp
View file @
f0a23d1b
...
@@ -65,12 +65,10 @@ struct cl_spawn_helper<R(Ts...)> {
...
@@ -65,12 +65,10 @@ struct cl_spawn_helper<R(Ts...)> {
}
// namespace detail
}
// namespace detail
/**
/// Creates a new actor facade for an OpenCL kernel that invokes
* Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
* the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
* @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
* `dims.empty()`, or `clCreateKernel` failed.
*/
template
<
class
Signature
>
template
<
class
Signature
>
actor
spawn_cl
(
const
opencl
::
program
&
prog
,
actor
spawn_cl
(
const
opencl
::
program
&
prog
,
const
char
*
fname
,
const
char
*
fname
,
...
@@ -82,13 +80,11 @@ actor spawn_cl(const opencl::program& prog,
...
@@ -82,13 +80,11 @@ actor spawn_cl(const opencl::program& prog,
return
f
(
prog
,
fname
,
dims
,
offset
,
local_dims
,
result_size
);
return
f
(
prog
,
fname
,
dims
,
offset
,
local_dims
,
result_size
);
}
}
/**
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
* Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
* that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
* @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
* <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
* occured, or @p clCreateKernel failed.
*/
template
<
class
Signature
>
template
<
class
Signature
>
actor
spawn_cl
(
const
char
*
source
,
actor
spawn_cl
(
const
char
*
source
,
const
char
*
fname
,
const
char
*
fname
,
...
@@ -100,12 +96,10 @@ actor spawn_cl(const char* source,
...
@@ -100,12 +96,10 @@ actor spawn_cl(const char* source,
dims
,
offset
,
local_dims
,
result_size
);
dims
,
offset
,
local_dims
,
result_size
);
}
}
/**
/// Creates a new actor facade for an OpenCL kernel that invokes
* Creates a new actor facade for an OpenCL kernel that invokes
/// the function named `fname` from `prog`.
* the function named `fname` from `prog`.
/// @throws std::runtime_error if more than three dimensions are set,
* @throws std::runtime_error if more than three dimensions are set,
/// `dims.empty()`, or `clCreateKernel` failed.
* `dims.empty()`, or `clCreateKernel` failed.
*/
template
<
class
Signature
,
class
Fun
>
template
<
class
Signature
,
class
Fun
>
actor
spawn_cl
(
const
opencl
::
program
&
prog
,
actor
spawn_cl
(
const
opencl
::
program
&
prog
,
const
char
*
fname
,
const
char
*
fname
,
...
@@ -120,13 +114,11 @@ actor spawn_cl(const opencl::program& prog,
...
@@ -120,13 +114,11 @@ actor spawn_cl(const opencl::program& prog,
std
::
move
(
map_args
),
std
::
move
(
map_result
));
std
::
move
(
map_args
),
std
::
move
(
map_result
));
}
}
/**
/// Compiles `source` and creates a new actor facade for an OpenCL kernel
* Compiles `source` and creates a new actor facade for an OpenCL kernel
/// that invokes the function named `fname`.
* that invokes the function named `fname`.
/// @throws std::runtime_error if more than three dimensions are set,
* @throws std::runtime_error if more than three dimensions are set,
/// <tt>dims.empty()</tt>, a compilation error
* <tt>dims.empty()</tt>, a compilation error
/// occured, or @p clCreateKernel failed.
* occured, or @p clCreateKernel failed.
*/
template
<
class
Signature
,
class
Fun
>
template
<
class
Signature
,
class
Fun
>
actor
spawn_cl
(
const
char
*
source
,
actor
spawn_cl
(
const
char
*
source
,
const
char
*
fname
,
const
char
*
fname
,
...
...
libcaf_opencl/examples/proper_matrix.cpp
View file @
f0a23d1b
...
@@ -59,7 +59,7 @@ constexpr const char* kernel_source = R"__(
...
@@ -59,7 +59,7 @@ constexpr const char* kernel_source = R"__(
template
<
size_t
Size
>
template
<
size_t
Size
>
class
square_matrix
{
class
square_matrix
{
public:
public:
static
constexpr
size_t
num_elements
=
Size
*
Size
;
static
constexpr
size_t
num_elements
=
Size
*
Size
;
...
@@ -68,38 +68,38 @@ class square_matrix {
...
@@ -68,38 +68,38 @@ class square_matrix {
square_matrix
&
operator
=
(
square_matrix
&&
)
=
default
;
square_matrix
&
operator
=
(
square_matrix
&&
)
=
default
;
square_matrix
&
operator
=
(
const
square_matrix
&
)
=
default
;
square_matrix
&
operator
=
(
const
square_matrix
&
)
=
default
;
square_matrix
()
:
m_data
(
num_elements
)
{
}
square_matrix
()
:
data_
(
num_elements
)
{
}
explicit
square_matrix
(
fvec
d
)
:
m_data
(
move
(
d
))
{
explicit
square_matrix
(
fvec
d
)
:
data_
(
move
(
d
))
{
assert
(
m_data
.
size
()
==
num_elements
);
assert
(
data_
.
size
()
==
num_elements
);
}
}
inline
float
&
operator
()(
size_t
column
,
size_t
row
)
{
inline
float
&
operator
()(
size_t
column
,
size_t
row
)
{
return
m_data
[
column
+
row
*
Size
];
return
data_
[
column
+
row
*
Size
];
}
}
inline
const
float
&
operator
()(
size_t
column
,
size_t
row
)
const
{
inline
const
float
&
operator
()(
size_t
column
,
size_t
row
)
const
{
return
m_data
[
column
+
row
*
Size
];
return
data_
[
column
+
row
*
Size
];
}
}
inline
void
iota_fill
()
{
inline
void
iota_fill
()
{
iota
(
m_data
.
begin
(),
m_data
.
end
(),
0
);
iota
(
data_
.
begin
(),
data_
.
end
(),
0
);
}
}
typedef
typename
fvec
::
const_iterator
const_iterator
;
typedef
typename
fvec
::
const_iterator
const_iterator
;
const_iterator
begin
()
const
{
return
m_data
.
begin
();
}
const_iterator
begin
()
const
{
return
data_
.
begin
();
}
const_iterator
end
()
const
{
return
m_data
.
end
();
}
const_iterator
end
()
const
{
return
data_
.
end
();
}
fvec
&
data
()
{
return
m_data
;
}
fvec
&
data
()
{
return
data_
;
}
const
fvec
&
data
()
const
{
return
m_data
;
}
const
fvec
&
data
()
const
{
return
data_
;
}
private:
private:
fvec
m_data
;
fvec
data_
;
};
};
...
@@ -126,7 +126,7 @@ inline bool operator==(const square_matrix<Size>& lhs,
...
@@ -126,7 +126,7 @@ inline bool operator==(const square_matrix<Size>& lhs,
template
<
size_t
Size
>
template
<
size_t
Size
>
inline
bool
operator
!=
(
const
square_matrix
<
Size
>&
lhs
,
inline
bool
operator
!=
(
const
square_matrix
<
Size
>&
lhs
,
const
square_matrix
<
Size
>&
rhs
)
{
const
square_matrix
<
Size
>&
rhs
)
{
return
!
(
lhs
==
rhs
);
return
!
(
lhs
==
rhs
);
}
}
using
matrix_type
=
square_matrix
<
matrix_size
>
;
using
matrix_type
=
square_matrix
<
matrix_size
>
;
...
...
libcaf_opencl/src/opencl_metainfo.cpp
View file @
f0a23d1b
...
@@ -32,7 +32,7 @@ opencl_metainfo* opencl_metainfo::instance() {
...
@@ -32,7 +32,7 @@ opencl_metainfo* opencl_metainfo::instance() {
}
}
const
std
::
vector
<
device_info
>
opencl_metainfo
::
get_devices
()
const
{
const
std
::
vector
<
device_info
>
opencl_metainfo
::
get_devices
()
const
{
return
m_devices
;
return
devices_
;
}
}
void
opencl_metainfo
::
initialize
()
{
void
opencl_metainfo
::
initialize
()
{
...
@@ -65,7 +65,7 @@ void opencl_metainfo::initialize() {
...
@@ -65,7 +65,7 @@ void opencl_metainfo::initialize() {
auto
lift
=
[](
cl_device_id
ptr
)
{
return
device_ptr
{
ptr
,
false
};
};
auto
lift
=
[](
cl_device_id
ptr
)
{
return
device_ptr
{
ptr
,
false
};
};
std
::
transform
(
ds
.
begin
(),
ds
.
end
(),
devices
.
begin
(),
lift
);
std
::
transform
(
ds
.
begin
(),
ds
.
end
(),
devices
.
begin
(),
lift
);
// create a context
// create a context
m_context
.
reset
(
v2get
(
CAF_CLF
(
clCreateContext
),
nullptr
,
num_devs
,
ds
.
data
(),
context_
.
reset
(
v2get
(
CAF_CLF
(
clCreateContext
),
nullptr
,
num_devs
,
ds
.
data
(),
pfn_notify
,
nullptr
),
pfn_notify
,
nullptr
),
false
);
false
);
for
(
auto
&
device
:
devices
)
{
for
(
auto
&
device
:
devices
)
{
...
@@ -73,7 +73,7 @@ void opencl_metainfo::initialize() {
...
@@ -73,7 +73,7 @@ void opencl_metainfo::initialize() {
command_queue_ptr
cmd_queue
;
command_queue_ptr
cmd_queue
;
try
{
try
{
cmd_queue
.
reset
(
v2get
(
CAF_CLF
(
clCreateCommandQueue
),
cmd_queue
.
reset
(
v2get
(
CAF_CLF
(
clCreateCommandQueue
),
m_context
.
get
(),
device
.
get
(),
context_
.
get
(),
device
.
get
(),
unsigned
{
CL_QUEUE_PROFILING_ENABLE
}),
unsigned
{
CL_QUEUE_PROFILING_ENABLE
}),
false
);
false
);
}
}
...
@@ -90,11 +90,11 @@ void opencl_metainfo::initialize() {
...
@@ -90,11 +90,11 @@ void opencl_metainfo::initialize() {
unsigned
{
CL_DEVICE_MAX_WORK_ITEM_SIZES
},
unsigned
{
CL_DEVICE_MAX_WORK_ITEM_SIZES
},
sizeof
(
size_t
)
*
max_wid
,
sizeof
(
size_t
)
*
max_wid
,
max_wi_per_dim
.
data
());
max_wi_per_dim
.
data
());
m_devices
.
push_back
(
device_info
{
std
::
move
(
device
),
std
::
move
(
cmd_queue
),
devices_
.
push_back
(
device_info
{
std
::
move
(
device
),
std
::
move
(
cmd_queue
),
max_wgs
,
max_wid
,
max_wi_per_dim
});
max_wgs
,
max_wid
,
max_wi_per_dim
});
}
}
}
}
if
(
m_devices
.
empty
())
{
if
(
devices_
.
empty
())
{
std
::
string
errstr
=
"could not create a command queue for any device"
;
std
::
string
errstr
=
"could not create a command queue for any device"
;
CAF_LOG_ERROR
(
errstr
);
CAF_LOG_ERROR
(
errstr
);
throw
std
::
runtime_error
(
std
::
move
(
errstr
));
throw
std
::
runtime_error
(
std
::
move
(
errstr
));
...
...
libcaf_opencl/src/program.cpp
View file @
f0a23d1b
...
@@ -35,15 +35,15 @@ namespace opencl {
...
@@ -35,15 +35,15 @@ namespace opencl {
program
::
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program
::
program
(
context_ptr
context
,
command_queue_ptr
queue
,
program_ptr
program
)
program_ptr
program
)
:
m_context
(
move
(
context
)),
:
context_
(
move
(
context
)),
m_program
(
move
(
program
)),
program_
(
move
(
program
)),
m_queue
(
move
(
queue
))
{}
queue_
(
move
(
queue
))
{}
program
program
::
create
(
const
char
*
kernel_source
,
const
char
*
options
,
program
program
::
create
(
const
char
*
kernel_source
,
const
char
*
options
,
uint32_t
device_id
)
{
uint32_t
device_id
)
{
auto
metainfo
=
opencl_metainfo
::
instance
();
auto
metainfo
=
opencl_metainfo
::
instance
();
auto
devices
=
metainfo
->
get_devices
();
auto
devices
=
metainfo
->
get_devices
();
auto
context
=
metainfo
->
m_context
;
auto
context
=
metainfo
->
context_
;
if
(
devices
.
size
()
<=
device_id
)
{
if
(
devices
.
size
()
<=
device_id
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"Device id "
<<
device_id
oss
<<
"Device id "
<<
device_id
...
@@ -59,7 +59,7 @@ program program::create(const char* kernel_source, const char* options,
...
@@ -59,7 +59,7 @@ program program::create(const char* kernel_source, const char* options,
cl_uint
{
1
},
&
kernel_source
,
&
kernel_source_length
);
cl_uint
{
1
},
&
kernel_source
,
&
kernel_source_length
);
pptr
.
reset
(
rawptr
,
false
);
pptr
.
reset
(
rawptr
,
false
);
// build programm from program object
// build programm from program object
auto
dev_tmp
=
devices
[
device_id
].
m_device
.
get
();
auto
dev_tmp
=
devices
[
device_id
].
device_
.
get
();
cl_int
err
=
clBuildProgram
(
pptr
.
get
(),
1
,
&
dev_tmp
,
cl_int
err
=
clBuildProgram
(
pptr
.
get
(),
1
,
&
dev_tmp
,
options
,
nullptr
,
nullptr
);
options
,
nullptr
,
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
...
@@ -91,7 +91,7 @@ program program::create(const char* kernel_source, const char* options,
...
@@ -91,7 +91,7 @@ program program::create(const char* kernel_source, const char* options,
#endif
#endif
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
return
{
context
,
devices
[
device_id
].
m_cmd_queue
,
pptr
};
return
{
context
,
devices
[
device_id
].
cmd_queue_
,
pptr
};
}
}
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/test/opencl.cpp
View file @
f0a23d1b
...
@@ -96,11 +96,11 @@ constexpr const char* kernel_source_const = R"__(
...
@@ -96,11 +96,11 @@ constexpr const char* kernel_source_const = R"__(
template
<
size_t
Size
>
template
<
size_t
Size
>
class
square_matrix
{
class
square_matrix
{
public:
public:
static
constexpr
size_t
num_elements
=
Size
*
Size
;
static
constexpr
size_t
num_elements
=
Size
*
Size
;
static
void
announce
()
{
static
void
announce
()
{
caf
::
announce
<
square_matrix
>
(
"square_matrix"
,
&
square_matrix
::
m_data
);
caf
::
announce
<
square_matrix
>
(
"square_matrix"
,
&
square_matrix
::
data_
);
}
}
square_matrix
(
square_matrix
&&
)
=
default
;
square_matrix
(
square_matrix
&&
)
=
default
;
...
@@ -108,46 +108,46 @@ class square_matrix {
...
@@ -108,46 +108,46 @@ class square_matrix {
square_matrix
&
operator
=
(
square_matrix
&&
)
=
default
;
square_matrix
&
operator
=
(
square_matrix
&&
)
=
default
;
square_matrix
&
operator
=
(
const
square_matrix
&
)
=
default
;
square_matrix
&
operator
=
(
const
square_matrix
&
)
=
default
;
square_matrix
()
:
m_data
(
num_elements
)
{
square_matrix
()
:
data_
(
num_elements
)
{
// nop
// nop
}
}
explicit
square_matrix
(
ivec
d
)
:
m_data
(
move
(
d
))
{
explicit
square_matrix
(
ivec
d
)
:
data_
(
move
(
d
))
{
assert
(
m_data
.
size
()
==
num_elements
);
assert
(
data_
.
size
()
==
num_elements
);
}
}
float
&
operator
()(
size_t
column
,
size_t
row
)
{
float
&
operator
()(
size_t
column
,
size_t
row
)
{
return
m_data
[
column
+
row
*
Size
];
return
data_
[
column
+
row
*
Size
];
}
}
const
float
&
operator
()(
size_t
column
,
size_t
row
)
const
{
const
float
&
operator
()(
size_t
column
,
size_t
row
)
const
{
return
m_data
[
column
+
row
*
Size
];
return
data_
[
column
+
row
*
Size
];
}
}
typedef
typename
ivec
::
const_iterator
const_iterator
;
typedef
typename
ivec
::
const_iterator
const_iterator
;
const_iterator
begin
()
const
{
const_iterator
begin
()
const
{
return
m_data
.
begin
();
return
data_
.
begin
();
}
}
const_iterator
end
()
const
{
const_iterator
end
()
const
{
return
m_data
.
end
();
return
data_
.
end
();
}
}
ivec
&
data
()
{
ivec
&
data
()
{
return
m_data
;
return
data_
;
}
}
const
ivec
&
data
()
const
{
const
ivec
&
data
()
const
{
return
m_data
;
return
data_
;
}
}
void
data
(
ivec
new_data
)
{
void
data
(
ivec
new_data
)
{
m_data
=
std
::
move
(
new_data
);
data_
=
std
::
move
(
new_data
);
}
}
private:
private:
ivec
m_data
;
ivec
data_
;
};
};
...
@@ -175,7 +175,7 @@ bool operator==(const square_matrix<Size>& lhs,
...
@@ -175,7 +175,7 @@ bool operator==(const square_matrix<Size>& lhs,
template
<
size_t
Size
>
template
<
size_t
Size
>
bool
operator
!=
(
const
square_matrix
<
Size
>&
lhs
,
bool
operator
!=
(
const
square_matrix
<
Size
>&
lhs
,
const
square_matrix
<
Size
>&
rhs
)
{
const
square_matrix
<
Size
>&
rhs
)
{
return
!
(
lhs
==
rhs
);
return
!
(
lhs
==
rhs
);
}
}
using
matrix_type
=
square_matrix
<
matrix_size
>
;
using
matrix_type
=
square_matrix
<
matrix_size
>
;
...
...
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