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
86e60176
Commit
86e60176
authored
Jul 26, 2014
by
Marian Triebe
Browse files
Options
Browse Files
Download
Email Patches
Plain Diff
Apply new coding style
parent
f2a33b6a
Changes
11
Show whitespace changes
Inline
Side-by-side
Showing
11 changed files
with
650 additions
and
775 deletions
+650
-775
libcaf_opencl/caf/opencl/actor_facade.hpp
libcaf_opencl/caf/opencl/actor_facade.hpp
+144
-162
libcaf_opencl/caf/opencl/command.hpp
libcaf_opencl/caf/opencl/command.hpp
+82
-98
libcaf_opencl/caf/opencl/device_info.hpp
libcaf_opencl/caf/opencl/device_info.hpp
+21
-28
libcaf_opencl/caf/opencl/global.hpp
libcaf_opencl/caf/opencl/global.hpp
+4
-4
libcaf_opencl/caf/opencl/opencl_metainfo.hpp
libcaf_opencl/caf/opencl/opencl_metainfo.hpp
+12
-15
libcaf_opencl/caf/opencl/program.hpp
libcaf_opencl/caf/opencl/program.hpp
+14
-16
libcaf_opencl/caf/opencl/smart_ptr.hpp
libcaf_opencl/caf/opencl/smart_ptr.hpp
+51
-53
libcaf_opencl/caf/opencl/spawn_cl.hpp
libcaf_opencl/caf/opencl/spawn_cl.hpp
+81
-114
libcaf_opencl/src/global.cpp
libcaf_opencl/src/global.cpp
+53
-53
libcaf_opencl/src/opencl_metainfo.cpp
libcaf_opencl/src/opencl_metainfo.cpp
+124
-158
libcaf_opencl/src/program.cpp
libcaf_opencl/src/program.cpp
+64
-74
No files found.
libcaf_opencl/caf/opencl/actor_facade.hpp
View file @
86e60176
...
@@ -55,7 +55,6 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -55,7 +55,6 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
friend
class
command
<
actor_facade
,
Ret
>
;
friend
class
command
<
actor_facade
,
Ret
>
;
public:
public:
typedef
cow_tuple
<
typename
detail
::
rm_const_and_ref
<
Args
>::
type
...
>
typedef
cow_tuple
<
typename
detail
::
rm_const_and_ref
<
Args
>::
type
...
>
args_tuple
;
args_tuple
;
...
@@ -63,14 +62,13 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -63,14 +62,13 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
typedef
std
::
function
<
any_tuple
(
Ret
&
)
>
result_mapping
;
typedef
std
::
function
<
any_tuple
(
Ret
&
)
>
result_mapping
;
static
intrusive_ptr
<
actor_facade
>
static
intrusive_ptr
<
actor_facade
>
create
(
const
program
&
prog
,
const
char
*
kernel_name
,
create
(
const
program
&
prog
,
const
char
*
kernel_name
,
arg_mapping
map_args
,
arg_mapping
map_args
,
result_mapping
map_result
,
result_mapping
map_result
,
const
dim_vec
&
global_dims
,
const
dim_vec
&
global_dims
,
const
dim_vec
&
offset
s
,
const
dim_vec
&
offsets
,
const
dim_vec
&
local_dim
s
,
const
dim_vec
&
local_dims
,
size_t
result_size
)
{
size_t
result_size
)
{
if
(
global_dims
.
empty
())
{
if
(
global_dims
.
empty
())
{
auto
str
=
"OpenCL kernel needs at least 1 global dimension."
;
auto
str
=
"OpenCL kernel needs at least 1 global dimension."
;
CPPA_LOGM_ERROR
(
detail
::
demangle
(
typeid
(
actor_facade
)).
c_str
(),
CPPA_LOGM_ERROR
(
detail
::
demangle
(
typeid
(
actor_facade
)).
c_str
(),
str
);
str
);
throw
std
::
runtime_error
(
str
);
throw
std
::
runtime_error
(
str
);
}
}
auto
check_vec
=
[
&
](
const
dim_vec
&
vec
,
const
char
*
name
)
{
auto
check_vec
=
[
&
](
const
dim_vec
&
vec
,
const
char
*
name
)
{
...
@@ -78,33 +76,28 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -78,33 +76,28 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
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"
;
CPPA_LOGM_ERROR
(
detail
::
demangle
<
actor_facade
>
().
c_str
(),
CPPA_LOGM_ERROR
(
detail
::
demangle
<
actor_facade
>
().
c_str
(),
oss
.
str
());
oss
.
str
());
throw
std
::
runtime_error
(
oss
.
str
());
throw
std
::
runtime_error
(
oss
.
str
());
}
}
};
};
check_vec
(
offsets
,
"offsets"
);
check_vec
(
offsets
,
"offsets"
);
check_vec
(
local_dims
,
"local dimensions"
);
check_vec
(
local_dims
,
"local dimensions"
);
cl_int
err
{
0
};
cl_int
err
{
0
};
kernel_ptr
kernel
;
kernel_ptr
kernel
;
kernel
.
adopt
(
clCreateKernel
(
prog
.
m_program
.
get
(),
kernel_name
,
&
err
));
kernel
.
adopt
(
clCreateKernel
(
prog
.
m_program
.
get
(),
kernel_name
,
&
err
));
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
std
::
ostringstream
oss
;
std
::
ostringstream
oss
;
oss
<<
"clCreateKernel: "
<<
get_opencl_error
(
err
);
oss
<<
"clCreateKernel: "
<<
get_opencl_error
(
err
);
CPPA_LOGM_ERROR
(
detail
::
demangle
<
actor_facade
>
().
c_str
(),
CPPA_LOGM_ERROR
(
detail
::
demangle
<
actor_facade
>
().
c_str
(),
oss
.
str
());
oss
.
str
());
throw
std
::
runtime_error
(
oss
.
str
());
throw
std
::
runtime_error
(
oss
.
str
());
}
}
if
(
result_size
==
0
)
{
if
(
result_size
==
0
)
{
result_size
=
std
::
accumulate
(
global_dims
.
begin
(),
result_size
=
std
::
accumulate
(
global_dims
.
begin
(),
global_dims
.
end
(),
1
,
global_dims
.
end
(),
1
,
std
::
multiplies
<
size_t
>
{});
std
::
multiplies
<
size_t
>
{});
}
}
return
new
actor_facade
<
Ret
(
Args
...)
>
{
return
new
actor_facade
<
Ret
(
Args
...)
>
{
prog
,
kernel
,
global_dims
,
offsets
,
prog
,
kernel
,
global_dims
,
offsets
,
local_dims
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
result_size
local_dims
,
std
::
move
(
map_args
),
std
::
move
(
map_result
),
result_size
};
};
}
}
void
enqueue
(
msg_hdr_cref
hdr
,
any_tuple
msg
,
execution_unit
*
)
override
{
void
enqueue
(
msg_hdr_cref
hdr
,
any_tuple
msg
,
execution_unit
*
)
override
{
...
@@ -114,25 +107,23 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -114,25 +107,23 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
}
}
private:
private:
using
evnt_vec
=
std
::
vector
<
cl_event
>
;
using
evnt_vec
=
std
::
vector
<
cl_event
>
;
using
args_vec
=
std
::
vector
<
mem_ptr
>
;
using
args_vec
=
std
::
vector
<
mem_ptr
>
;
actor_facade
(
const
program
&
prog
,
kernel_ptr
kernel
,
actor_facade
(
const
program
&
prog
,
kernel_ptr
kernel
,
const
dim_vec
&
global_dimension
s
,
const
dim_vec
&
global_dimensions
,
const
dim_vec
&
global_offset
s
,
const
dim_vec
&
global_offset
s
,
const
dim_vec
&
local_dimensions
,
arg_mapping
map_arg
s
,
const
dim_vec
&
local_dimensions
,
result_mapping
map_result
,
size_t
result_size
)
arg_mapping
map_args
,
result_mapping
map_result
,
:
m_kernel
(
kernel
)
size_t
result_size
)
,
m_program
(
prog
.
m_program
)
:
m_kernel
(
kernel
)
,
m_program
(
prog
.
m_program
)
,
m_context
(
prog
.
m_context
)
,
m_
context
(
prog
.
m_context
)
,
m_
queue
(
prog
.
m_queue
)
,
m_queue
(
prog
.
m_queue
)
,
m_global_dimensions
(
global_dimensions
)
,
m_global_dimensions
(
global_dimensions
)
,
m_global_offsets
(
global_offsets
)
,
m_global_offsets
(
global_offsets
)
,
m_local_dimensions
(
local_dimensions
)
,
m_local_dimensions
(
local_dimensions
)
,
m_map_args
(
std
::
move
(
map_args
))
,
m_map_args
(
std
::
move
(
map_args
))
,
m_map_result
(
std
::
move
(
map_result
))
,
m_map_result
(
std
::
move
(
map_result
))
,
m_result_size
(
result_size
)
,
m_result_size
(
result_size
)
{
{
CPPA_LOG_TRACE
(
"id: "
<<
this
->
id
());
CPPA_LOG_TRACE
(
"id: "
<<
this
->
id
());
}
}
...
@@ -147,17 +138,14 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -147,17 +138,14 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
add_arguments_to_kernel
<
Ret
>
(
events
,
arguments
,
m_result_size
,
add_arguments_to_kernel
<
Ret
>
(
events
,
arguments
,
m_result_size
,
get_ref
<
Is
>
(
*
opt
)...);
get_ref
<
Is
>
(
*
opt
)...);
auto
cmd
=
make_counted
<
command
<
actor_facade
,
Ret
>>
(
auto
cmd
=
make_counted
<
command
<
actor_facade
,
Ret
>>
(
handle
,
this
,
handle
,
this
,
std
::
move
(
events
),
std
::
move
(
arguments
),
m_result_size
,
std
::
move
(
events
),
std
::
move
(
arguments
),
*
opt
);
m_result_size
,
*
opt
);
cmd
->
enqueue
();
cmd
->
enqueue
();
}
else
{
}
else
{
CPPA_LOGMF
(
CPPA_ERROR
,
"actor_facade::enqueue() tuple_cast failed."
);
CPPA_LOGMF
(
CPPA_ERROR
,
"actor_facade::enqueue() tuple_cast failed."
);
}
}
}
}
kernel_ptr
m_kernel
;
kernel_ptr
m_kernel
;
program_ptr
m_program
;
program_ptr
m_program
;
context_ptr
m_context
;
context_ptr
m_context
;
...
@@ -173,7 +161,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -173,7 +161,7 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
cl_int
err
{
0
};
cl_int
err
{
0
};
// rotate left (output buffer to the end)
// rotate left (output buffer to the end)
rotate
(
begin
(
arguments
),
begin
(
arguments
)
+
1
,
end
(
arguments
));
rotate
(
begin
(
arguments
),
begin
(
arguments
)
+
1
,
end
(
arguments
));
for
(
size_t
i
=
0
;
i
<
arguments
.
size
();
++
i
)
{
for
(
size_t
i
=
0
;
i
<
arguments
.
size
();
++
i
)
{
err
=
clSetKernelArg
(
m_kernel
.
get
(),
i
,
sizeof
(
cl_mem
),
err
=
clSetKernelArg
(
m_kernel
.
get
(),
i
,
sizeof
(
cl_mem
),
static_cast
<
void
*>
(
&
arguments
[
i
]));
static_cast
<
void
*>
(
&
arguments
[
i
]));
CPPA_LOG_ERROR_IF
(
err
!=
CL_SUCCESS
,
CPPA_LOG_ERROR_IF
(
err
!=
CL_SUCCESS
,
...
@@ -182,27 +170,22 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -182,27 +170,22 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
clFlush
(
m_queue
.
get
());
clFlush
(
m_queue
.
get
());
}
}
template
<
typename
T0
,
typename
...
Ts
>
template
<
typename
T0
,
typename
...
Ts
>
void
add_arguments_to_kernel_rec
(
evnt_vec
&
events
,
args_vec
&
arguments
,
void
add_arguments_to_kernel_rec
(
evnt_vec
&
events
,
args_vec
&
arguments
,
T0
&
arg0
,
Ts
&
...
args
)
{
T0
&
arg0
,
Ts
&
...
args
)
{
cl_int
err
{
0
};
cl_int
err
{
0
};
size_t
buffer_size
=
sizeof
(
typename
T0
::
value_type
)
*
arg0
.
size
();
size_t
buffer_size
=
sizeof
(
typename
T0
::
value_type
)
*
arg0
.
size
();
auto
buffer
=
clCreateBuffer
(
m_context
.
get
(),
auto
buffer
=
clCreateBuffer
(
m_context
.
get
(),
CL_MEM_READ_ONLY
,
buffer_size
,
CL_MEM_READ_ONLY
,
nullptr
,
&
err
);
buffer_size
,
nullptr
,
&
err
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_ERROR
,
"clCreateBuffer: "
<<
get_opencl_error
(
err
));
CPPA_LOGMF
(
CPPA_ERROR
,
"clCreateBuffer: "
<<
get_opencl_error
(
err
));
return
;
return
;
}
}
cl_event
event
;
cl_event
event
;
err
=
clEnqueueWriteBuffer
(
m_queue
.
get
(),
buffer
,
CL_FALSE
,
0
,
err
=
clEnqueueWriteBuffer
(
m_queue
.
get
(),
buffer
,
CL_FALSE
,
0
,
buffer_size
,
buffer_size
,
arg0
.
data
(),
arg0
.
data
(),
0
,
nullptr
,
&
event
);
0
,
nullptr
,
&
event
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_ERROR
,
"clEnqueueWriteBuffer: "
CPPA_LOGMF
(
CPPA_ERROR
,
"clEnqueueWriteBuffer: "
<<
get_opencl_error
(
err
));
<<
get_opencl_error
(
err
));
return
;
return
;
}
}
events
.
push_back
(
std
::
move
(
event
));
events
.
push_back
(
std
::
move
(
event
));
...
@@ -212,14 +195,14 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -212,14 +195,14 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
add_arguments_to_kernel_rec
(
events
,
arguments
,
args
...);
add_arguments_to_kernel_rec
(
events
,
arguments
,
args
...);
}
}
template
<
typename
R
,
typename
...
Ts
>
template
<
typename
R
,
typename
...
Ts
>
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
,
Ts
&&
...
args
)
{
size_t
ret_size
,
Ts
&&
...
args
)
{
arguments
.
clear
();
arguments
.
clear
();
cl_int
err
{
0
};
cl_int
err
{
0
};
auto
buf
=
clCreateBuffer
(
m_context
.
get
(),
CL_MEM_WRITE_ONLY
,
auto
buf
=
sizeof
(
typename
R
::
value_type
)
*
ret_size
,
clCreateBuffer
(
m_context
.
get
(),
CL_MEM_WRITE_ONLY
,
nullptr
,
&
err
);
sizeof
(
typename
R
::
value_type
)
*
ret_size
,
nullptr
,
&
err
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_ERROR
,
"clCreateBuffer: "
<<
get_opencl_error
(
err
));
CPPA_LOGMF
(
CPPA_ERROR
,
"clCreateBuffer: "
<<
get_opencl_error
(
err
));
return
;
return
;
...
@@ -229,7 +212,6 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
...
@@ -229,7 +212,6 @@ class actor_facade<Ret(Args...)> : public abstract_actor {
arguments
.
push_back
(
tmp
);
arguments
.
push_back
(
tmp
);
add_arguments_to_kernel_rec
(
events
,
arguments
,
std
::
forward
<
Ts
>
(
args
)...);
add_arguments_to_kernel_rec
(
events
,
arguments
,
std
::
forward
<
Ts
>
(
args
)...);
}
}
};
};
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/caf/opencl/command.hpp
View file @
86e60176
...
@@ -35,17 +35,13 @@
...
@@ -35,17 +35,13 @@
namespace
caf
{
namespace
caf
{
namespace
opencl
{
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
,
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
mem_ptr
>
arguments
,
intrusive_ptr
<
T
>
actor_facade
,
size_t
result_size
,
any_tuple
msg
)
std
::
vector
<
cl_event
>
events
,
std
::
vector
<
mem_ptr
>
arguments
,
size_t
result_size
,
any_tuple
msg
)
:
m_result_size
(
result_size
)
:
m_result_size
(
result_size
)
,
m_handle
(
handle
)
,
m_handle
(
handle
)
,
m_actor_facade
(
actor_facade
)
,
m_actor_facade
(
actor_facade
)
...
@@ -53,20 +49,19 @@ class command : public ref_counted {
...
@@ -53,20 +49,19 @@ class command : public ref_counted {
,
m_events
(
std
::
move
(
events
))
,
m_events
(
std
::
move
(
events
))
,
m_arguments
(
std
::
move
(
arguments
))
,
m_arguments
(
std
::
move
(
arguments
))
,
m_result
(
m_result_size
)
,
m_result
(
m_result_size
)
,
m_msg
(
msg
)
{
}
,
m_msg
(
msg
)
{
}
~
command
()
{
~
command
()
{
cl_int
err
{
0
};
cl_int
err
{
0
};
for
(
auto
&
e
:
m_events
)
{
for
(
auto
&
e
:
m_events
)
{
err
=
clReleaseEvent
(
e
);
err
=
clReleaseEvent
(
e
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_ERROR
,
"clReleaseEvent: "
CPPA_LOGMF
(
CPPA_ERROR
,
"clReleaseEvent: "
<<
get_opencl_error
(
err
));
<<
get_opencl_error
(
err
));
}
}
}
}
}
}
void
enqueue
()
{
void
enqueue
()
{
CPPA_LOG_TRACE
(
"command::enqueue()"
);
CPPA_LOG_TRACE
(
"command::enqueue()"
);
this
->
ref
();
// reference held by the OpenCL comand queue
this
->
ref
();
// reference held by the OpenCL comand queue
cl_int
err
{
0
};
cl_int
err
{
0
};
...
@@ -74,40 +69,31 @@ class command : public ref_counted {
...
@@ -74,40 +69,31 @@ class command : public ref_counted {
auto
data_or_nullptr
=
[](
const
dim_vec
&
vec
)
{
auto
data_or_nullptr
=
[](
const
dim_vec
&
vec
)
{
return
vec
.
empty
()
?
nullptr
:
vec
.
data
();
return
vec
.
empty
()
?
nullptr
:
vec
.
data
();
};
};
err
=
clEnqueueNDRangeKernel
(
m_queue
.
get
(),
err
=
clEnqueueNDRangeKernel
(
m_actor_facade
->
m_kernel
.
get
(),
m_queue
.
get
(),
m_actor_facade
->
m_kernel
.
get
(),
m_actor_facade
->
m_global_dimensions
.
size
(),
m_actor_facade
->
m_global_dimensions
.
size
(),
data_or_nullptr
(
m_actor_facade
->
m_global_offsets
),
data_or_nullptr
(
m_actor_facade
->
m_global_offsets
),
data_or_nullptr
(
m_actor_facade
->
m_global_dimensions
),
data_or_nullptr
(
m_actor_facade
->
m_global_dimensions
),
data_or_nullptr
(
m_actor_facade
->
m_local_dimensions
),
data_or_nullptr
(
m_actor_facade
->
m_local_dimensions
),
m_events
.
size
(),
m_events
.
size
(),
(
m_events
.
empty
()
?
nullptr
:
m_events
.
data
()),
&
event_k
);
(
m_events
.
empty
()
?
nullptr
:
m_events
.
data
()),
&
event_k
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_ERROR
,
"clEnqueueNDRangeKernel: "
CPPA_LOGMF
(
CPPA_ERROR
,
<<
get_opencl_error
(
err
));
"clEnqueueNDRangeKernel: "
<<
get_opencl_error
(
err
));
this
->
deref
();
// or can anything actually happen?
this
->
deref
();
// or can anything actually happen?
return
;
return
;
}
}
else
{
else
{
cl_event
event_r
;
cl_event
event_r
;
err
=
clEnqueueReadBuffer
(
m_queue
.
get
(),
err
=
m_arguments
.
back
().
get
(),
clEnqueueReadBuffer
(
m_queue
.
get
(),
m_arguments
.
back
().
get
(),
CL_FALSE
,
CL_FALSE
,
0
,
sizeof
(
typename
R
::
value_type
)
*
m_result_size
,
0
,
m_result
.
data
(),
1
,
&
event_k
,
&
event_r
);
sizeof
(
typename
R
::
value_type
)
*
m_result_size
,
m_result
.
data
(),
1
,
&
event_k
,
&
event_r
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
throw
std
::
runtime_error
(
"clEnqueueReadBuffer: "
throw
std
::
runtime_error
(
"clEnqueueReadBuffer: "
+
+
get_opencl_error
(
err
));
get_opencl_error
(
err
));
this
->
deref
();
// failed to enqueue command
this
->
deref
();
// failed to enqueue command
return
;
return
;
}
}
err
=
clSetEventCallback
(
event_r
,
err
=
clSetEventCallback
(
event_r
,
CL_COMPLETE
,
CL_COMPLETE
,
[](
cl_event
,
cl_int
,
void
*
data
)
{
[](
cl_event
,
cl_int
,
void
*
data
)
{
auto
cmd
=
reinterpret_cast
<
command
*>
(
data
);
auto
cmd
=
reinterpret_cast
<
command
*>
(
data
);
cmd
->
handle_results
();
cmd
->
handle_results
();
...
@@ -115,8 +101,7 @@ class command : public ref_counted {
...
@@ -115,8 +101,7 @@ class command : public ref_counted {
},
},
this
);
this
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_ERROR
,
"clSetEventCallback: "
CPPA_LOGMF
(
CPPA_ERROR
,
"clSetEventCallback: "
<<
get_opencl_error
(
err
));
<<
get_opencl_error
(
err
));
this
->
deref
();
// callback is not set
this
->
deref
();
// callback is not set
return
;
return
;
}
}
...
@@ -131,7 +116,6 @@ class command : public ref_counted {
...
@@ -131,7 +116,6 @@ class command : public ref_counted {
}
}
private:
private:
int
m_result_size
;
int
m_result_size
;
response_promise
m_handle
;
response_promise
m_handle
;
intrusive_ptr
<
T
>
m_actor_facade
;
intrusive_ptr
<
T
>
m_actor_facade
;
...
@@ -139,9 +123,9 @@ class command : public ref_counted {
...
@@ -139,9 +123,9 @@ class command : public ref_counted {
std
::
vector
<
cl_event
>
m_events
;
std
::
vector
<
cl_event
>
m_events
;
std
::
vector
<
mem_ptr
>
m_arguments
;
std
::
vector
<
mem_ptr
>
m_arguments
;
R
m_result
;
R
m_result
;
any_tuple
m_msg
;
// required to keep the argument buffers alive (for
async copy)
any_tuple
m_msg
;
// required to keep the argument buffers alive (
async copy)
void
handle_results
()
{
void
handle_results
()
{
m_handle
.
deliver
(
m_actor_facade
->
m_map_result
(
m_result
));
m_handle
.
deliver
(
m_actor_facade
->
m_map_result
(
m_result
));
}
}
};
};
...
...
libcaf_opencl/caf/opencl/device_info.hpp
View file @
86e60176
...
@@ -32,24 +32,20 @@ class device_info {
...
@@ -32,24 +32,20 @@ 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
,
size_t
work_group_size
,
cl_uint
dimensons
,
command_queue_ptr
queue
,
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_max_work_group_size
(
work_group_size
)
,
m_max_dimensions
(
dimensons
)
,
m_max_dimensions
(
dimensons
)
,
m_max_work_items_per_dim
(
items_per_dimension
)
,
m_max_work_items_per_dim
(
items_per_dimension
)
,
m_device
(
device
)
,
m_device
(
device
)
,
m_cmd_queue
(
queue
)
{
}
,
m_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_max_work_group_size
;
cl_uint
m_max_dimensions
;
cl_uint
m_max_dimensions
;
dim_vec
m_max_work_items_per_dim
;
dim_vec
m_max_work_items_per_dim
;
...
@@ -65,9 +61,7 @@ inline size_t device_info::get_max_work_group_size() {
...
@@ -65,9 +61,7 @@ inline size_t device_info::get_max_work_group_size() {
return
m_max_work_group_size
;
return
m_max_work_group_size
;
}
}
inline
cl_uint
device_info
::
get_max_dimensions
()
{
inline
cl_uint
device_info
::
get_max_dimensions
()
{
return
m_max_dimensions
;
}
return
m_max_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_max_work_items_per_dim
;
...
@@ -76,5 +70,4 @@ inline dim_vec device_info::get_max_work_items_per_dim() {
...
@@ -76,5 +70,4 @@ inline dim_vec device_info::get_max_work_items_per_dim() {
}
// namespace opencl
}
// namespace opencl
}
// namespace caf
}
// namespace caf
#endif // CAF_OPENCL_DEVICE_INFO_HPP
#endif // CAF_OPENCL_DEVICE_INFO_HPP
libcaf_opencl/caf/opencl/global.hpp
View file @
86e60176
...
@@ -25,9 +25,9 @@
...
@@ -25,9 +25,9 @@
#include "caf/detail/limited_vector.hpp"
#include "caf/detail/limited_vector.hpp"
#if defined __APPLE__ || defined(MACOSX)
#if defined __APPLE__ || defined(MACOSX)
#include <OpenCL/opencl.h>
#include <OpenCL/opencl.h>
#else
#else
#include <CL/opencl.h>
#include <CL/opencl.h>
#endif
#endif
namespace
caf
{
namespace
caf
{
...
@@ -40,8 +40,8 @@ typedef detail::limited_vector<size_t, 3> dim_vec;
...
@@ -40,8 +40,8 @@ typedef detail::limited_vector<size_t, 3> dim_vec;
std
::
string
get_opencl_error
(
cl_int
err
);
std
::
string
get_opencl_error
(
cl_int
err
);
cl_int
clReleaseDeviceDummy
(
cl_device_id
);
cl_int
clReleaseDeviceDummy
(
cl_device_id
);
cl_int
clRetainDeviceDummy
(
cl_device_id
);
cl_int
clRetainDeviceDummy
(
cl_device_id
);
}
// namespace opencl
}
// namespace opencl
}
// namespace caf
}
// namespace caf
...
...
libcaf_opencl/caf/opencl/opencl_metainfo.hpp
View file @
86e60176
...
@@ -46,11 +46,9 @@ class opencl_metainfo {
...
@@ -46,11 +46,9 @@ class opencl_metainfo {
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
;
private:
private:
static
inline
opencl_metainfo
*
create_singleton
()
{
static
inline
opencl_metainfo
*
create_singleton
()
{
return
new
opencl_metainfo
;
return
new
opencl_metainfo
;
}
}
...
@@ -61,7 +59,6 @@ class opencl_metainfo {
...
@@ -61,7 +59,6 @@ class opencl_metainfo {
context_ptr
m_context
;
context_ptr
m_context
;
std
::
vector
<
device_info
>
m_devices
;
std
::
vector
<
device_info
>
m_devices
;
};
};
opencl_metainfo
*
get_opencl_metainfo
();
opencl_metainfo
*
get_opencl_metainfo
();
...
...
libcaf_opencl/caf/opencl/program.hpp
View file @
86e60176
...
@@ -28,7 +28,7 @@
...
@@ -28,7 +28,7 @@
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
template
<
typename
Signature
>
template
<
typename
Signature
>
class
actor_facade
;
class
actor_facade
;
/**
/**
...
@@ -36,26 +36,24 @@ class actor_facade;
...
@@ -36,26 +36,24 @@ class actor_facade;
*/
*/
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 cppa::opencl::program
* @brief Factory method, that creates a cppa::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
,
const
char
*
options
=
nullptr
,
uint32_t
device_id
=
0
);
static
program
create
(
const
char
*
kernel_source
,
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
m_context
;
program_ptr
m_program
;
program_ptr
m_program
;
command_queue_ptr
m_queue
;
command_queue_ptr
m_queue
;
};
};
}
// namespace opencl
}
// namespace opencl
...
...
libcaf_opencl/caf/opencl/smart_ptr.hpp
View file @
86e60176
...
@@ -27,7 +27,7 @@
...
@@ -27,7 +27,7 @@
namespace
caf
{
namespace
caf
{
namespace
opencl
{
namespace
opencl
{
template
<
typename
T
,
cl_int
(
*
ref
)(
T
),
cl_int
(
*
deref
)(
T
)>
template
<
typename
T
,
cl_int
(
*
ref
)(
T
),
cl_int
(
*
deref
)(
T
)>
class
smart_ptr
{
class
smart_ptr
{
typedef
typename
std
::
remove_pointer
<
T
>::
type
element_type
;
typedef
typename
std
::
remove_pointer
<
T
>::
type
element_type
;
...
@@ -37,22 +37,20 @@ class smart_ptr {
...
@@ -37,22 +37,20 @@ class smart_ptr {
typedef
const
element_type
*
const_pointer
;
typedef
const
element_type
*
const_pointer
;
typedef
const
element_type
&
const_reference
;
typedef
const
element_type
&
const_reference
;
public:
public:
smart_ptr
(
pointer
ptr
=
nullptr
)
:
m_ptr
(
ptr
)
{
smart_ptr
(
pointer
ptr
=
nullptr
)
:
m_ptr
(
ptr
)
{
if
(
m_ptr
)
ref
(
m_ptr
);
if
(
m_ptr
)
ref
(
m_ptr
);
}
}
~
smart_ptr
()
{
reset
();
}
~
smart_ptr
()
{
reset
();
}
smart_ptr
(
const
smart_ptr
&
other
)
:
m_ptr
(
other
.
m_ptr
)
{
smart_ptr
(
const
smart_ptr
&
other
)
:
m_ptr
(
other
.
m_ptr
)
{
if
(
m_ptr
)
ref
(
m_ptr
);
if
(
m_ptr
)
ref
(
m_ptr
);
}
}
smart_ptr
(
smart_ptr
&&
other
)
:
m_ptr
(
other
.
m_ptr
)
{
smart_ptr
(
smart_ptr
&&
other
)
:
m_ptr
(
other
.
m_ptr
)
{
other
.
m_ptr
=
nullptr
;
}
other
.
m_ptr
=
nullptr
;
}
smart_ptr
&
operator
=
(
pointer
ptr
)
{
smart_ptr
&
operator
=
(
pointer
ptr
)
{
reset
(
ptr
);
reset
(
ptr
);
...
@@ -71,9 +69,11 @@ class smart_ptr {
...
@@ -71,9 +69,11 @@ class smart_ptr {
}
}
inline
void
reset
(
pointer
ptr
=
nullptr
)
{
inline
void
reset
(
pointer
ptr
=
nullptr
)
{
if
(
m_ptr
)
deref
(
m_ptr
);
if
(
m_ptr
)
deref
(
m_ptr
);
m_ptr
=
ptr
;
m_ptr
=
ptr
;
if
(
ptr
)
ref
(
ptr
);
if
(
ptr
)
ref
(
ptr
);
}
}
// does not modify reference count of ptr
// does not modify reference count of ptr
...
@@ -93,9 +93,7 @@ class smart_ptr {
...
@@ -93,9 +93,7 @@ class smart_ptr {
inline
explicit
operator
bool
()
const
{
return
m_ptr
!=
nullptr
;
}
inline
explicit
operator
bool
()
const
{
return
m_ptr
!=
nullptr
;
}
private:
private:
pointer
m_ptr
;
pointer
m_ptr
;
};
};
typedef
smart_ptr
<
cl_mem
,
clRetainMemObject
,
clReleaseMemObject
>
mem_ptr
;
typedef
smart_ptr
<
cl_mem
,
clRetainMemObject
,
clReleaseMemObject
>
mem_ptr
;
...
@@ -104,9 +102,9 @@ typedef smart_ptr<cl_kernel, clRetainKernel, clReleaseKernel> kernel_ptr;
...
@@ -104,9 +102,9 @@ typedef smart_ptr<cl_kernel, clRetainKernel, clReleaseKernel> kernel_ptr;
typedef
smart_ptr
<
cl_context
,
clRetainContext
,
clReleaseContext
>
context_ptr
;
typedef
smart_ptr
<
cl_context
,
clRetainContext
,
clReleaseContext
>
context_ptr
;
typedef
smart_ptr
<
cl_program
,
clRetainProgram
,
clReleaseProgram
>
program_ptr
;
typedef
smart_ptr
<
cl_program
,
clRetainProgram
,
clReleaseProgram
>
program_ptr
;
typedef
smart_ptr
<
cl_device_id
,
clRetainDeviceDummy
,
clReleaseDeviceDummy
>
typedef
smart_ptr
<
cl_device_id
,
clRetainDeviceDummy
,
clReleaseDeviceDummy
>
device_ptr
;
device_ptr
;
typedef
smart_ptr
<
cl_command_queue
,
clRetainCommandQueue
,
clReleaseCommandQueue
>
typedef
smart_ptr
<
cl_command_queue
,
clRetainCommandQueue
,
clReleaseCommandQueue
>
command_queue_ptr
;
command_queue_ptr
;
}
// namespace opencl
}
// namespace opencl
}
// namespace caf
}
// namespace caf
...
...
libcaf_opencl/caf/opencl/spawn_cl.hpp
View file @
86e60176
...
@@ -38,64 +38,57 @@ namespace caf {
...
@@ -38,64 +38,57 @@ namespace caf {
namespace
detail
{
namespace
detail
{
// converts C arrays, i.e., pointers, to vectors
// converts C arrays, i.e., pointers, to vectors
template
<
typename
T
>
template
<
typename
T
>
struct
carr_to_vec
{
typedef
T
type
;
};
struct
carr_to_vec
{
typedef
T
type
;
};
template
<
typename
T
>
template
<
typename
T
>
struct
carr_to_vec
<
T
*>
{
typedef
std
::
vector
<
T
>
type
;
};
struct
carr_to_vec
<
T
*>
{
typedef
std
::
vector
<
T
>
type
;
};
template
<
typename
Signature
,
typename
SecondSignature
=
void
>
template
<
typename
Signature
,
typename
SecondSignature
=
void
>
struct
cl_spawn_helper
;
struct
cl_spawn_helper
;
template
<
typename
R
,
typename
...
Ts
>
template
<
typename
R
,
typename
...
Ts
>
struct
cl_spawn_helper
<
R
(
Ts
...),
void
>
{
struct
cl_spawn_helper
<
R
(
Ts
...),
void
>
{
using
result_type
=
typename
carr_to_vec
<
R
>::
type
;
using
result_type
=
typename
carr_to_vec
<
R
>::
type
;
using
impl
=
opencl
::
actor_facade
<
using
impl
=
opencl
::
actor_facade
<
result_type
(
typename
carr_to_vec
<
result_type
(
typename
carr_to_vec
<
typename
carr_to_vec
<
Ts
>::
type
>::
type
...)
>
;
typename
carr_to_vec
<
Ts
>::
type
>::
type
...)
>
;
using
map_arg_fun
=
typename
impl
::
arg_mapping
;
using
map_arg_fun
=
typename
impl
::
arg_mapping
;
using
map_res_fun
=
typename
impl
::
result_mapping
;
using
map_res_fun
=
typename
impl
::
result_mapping
;
template
<
typename
...
Us
>
template
<
typename
...
Us
>
actor
operator
()(
map_arg_fun
f0
,
actor
operator
()(
map_arg_fun
f0
,
map_res_fun
f1
,
const
opencl
::
program
&
p
,
map_res_fun
f1
,
const
char
*
fname
,
Us
&&
...
args
)
const
{
const
opencl
::
program
&
p
,
const
char
*
fname
,
Us
&&
...
args
)
const
{
using
std
::
move
;
using
std
::
move
;
using
std
::
forward
;
using
std
::
forward
;
return
impl
::
create
(
p
,
fname
,
move
(
f0
),
move
(
f1
),
forward
<
Us
>
(
args
)...);
return
impl
::
create
(
p
,
fname
,
move
(
f0
),
move
(
f1
),
forward
<
Us
>
(
args
)...);
}
}
template
<
typename
...
Us
>
template
<
typename
...
Us
>
actor
operator
()(
const
opencl
::
program
&
p
,
actor
operator
()(
const
opencl
::
program
&
p
,
const
char
*
fname
,
const
char
*
fname
,
Us
&&
...
args
)
const
{
Us
&&
...
args
)
const
{
using
std
::
move
;
using
std
::
move
;
using
std
::
forward
;
using
std
::
forward
;
map_arg_fun
f0
=
[]
(
any_tuple
msg
)
{
map_arg_fun
f0
=
[](
any_tuple
msg
)
{
return
tuple_cast
<
return
tuple_cast
<
typename
util
::
rm_const_and_ref
<
typename
util
::
rm_const_and_ref
<
typename
carr_to_vec
<
Ts
>::
type
>::
type
...
>
(
msg
);
typename
carr_to_vec
<
Ts
>::
type
>::
type
...
>
(
msg
);
};
};
map_res_fun
f1
=
[]
(
result_type
&
result
)
{
map_res_fun
f1
=
[]
(
result_type
&
result
)
{
return
make_any_tuple
(
move
(
result
));
return
make_any_tuple
(
move
(
result
));
};
};
return
impl
::
create
(
p
,
fname
,
move
(
f0
),
move
(
f1
),
forward
<
Us
>
(
args
)...);
return
impl
::
create
(
p
,
fname
,
move
(
f0
),
move
(
f1
),
forward
<
Us
>
(
args
)...);
}
}
};
};
template
<
typename
R
,
typename
...
Ts
>
template
<
typename
R
,
typename
...
Ts
>
struct
cl_spawn_helper
<
std
::
function
<
optional
<
cow_tuple
<
Ts
...
>>
(
any_tuple
)
>
,
struct
cl_spawn_helper
<
std
::
function
<
optional
<
cow_tuple
<
Ts
...
>>
(
any_tuple
)
>
,
std
::
function
<
any_tuple
(
R
&
)
>>
std
::
function
<
any_tuple
(
R
&
)
>>
:
cl_spawn_helper
<
R
(
Ts
...)
>
{
};
:
cl_spawn_helper
<
R
(
Ts
...)
>
{
};
}
// namespace detail
}
// namespace detail
...
@@ -106,13 +99,11 @@ struct cl_spawn_helper<std::function<optional<cow_tuple<Ts...>> (any_tuple)>,
...
@@ -106,13 +99,11 @@ struct cl_spawn_helper<std::function<optional<cow_tuple<Ts...>> (any_tuple)>,
* <tt>dims.empty()</tt>, or @p clCreateKernel
* <tt>dims.empty()</tt>, or @p clCreateKernel
* failed.
* failed.
*/
*/
template
<
typename
Signature
,
typename
...
Ts
>
template
<
typename
Signature
,
typename
...
Ts
>
inline
actor
spawn_cl
(
const
opencl
::
program
&
prog
,
inline
actor
const
char
*
fname
,
spawn_cl
(
const
opencl
::
program
&
prog
,
const
char
*
fname
,
const
opencl
::
dim_vec
&
dims
,
const
opencl
::
dim_vec
&
dims
,
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
local_dims
=
{},
size_t
result_size
=
0
)
{
const
opencl
::
dim_vec
&
local_dims
=
{},
size_t
result_size
=
0
)
{
using
std
::
move
;
using
std
::
move
;
detail
::
cl_spawn_helper
<
Signature
>
f
;
detail
::
cl_spawn_helper
<
Signature
>
f
;
return
f
(
prog
,
fname
,
dims
,
offset
,
local_dims
,
result_size
);
return
f
(
prog
,
fname
,
dims
,
offset
,
local_dims
,
result_size
);
...
@@ -125,20 +116,14 @@ inline actor spawn_cl(const opencl::program& prog,
...
@@ -125,20 +116,14 @@ inline actor spawn_cl(const opencl::program& prog,
* <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
<
typename
Signature
,
typename
...
Ts
>
template
<
typename
Signature
,
typename
...
Ts
>
inline
actor
spawn_cl
(
const
char
*
source
,
inline
actor
const
char
*
fname
,
spawn_cl
(
const
char
*
source
,
const
char
*
fname
,
const
opencl
::
dim_vec
&
dims
,
const
opencl
::
dim_vec
&
dims
,
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
local_dims
=
{},
const
opencl
::
dim_vec
&
local_dims
=
{},
size_t
result_size
=
0
)
{
size_t
result_size
=
0
)
{
using
std
::
move
;
using
std
::
move
;
return
spawn_cl
<
Signature
,
Ts
...
>
(
opencl
::
program
::
create
(
source
),
return
spawn_cl
<
Signature
,
Ts
...
>
(
opencl
::
program
::
create
(
source
),
fname
,
fname
,
dims
,
offset
,
local_dims
,
result_size
);
dims
,
offset
,
local_dims
,
result_size
);
}
}
/**
/**
...
@@ -150,27 +135,18 @@ inline actor spawn_cl(const char* source,
...
@@ -150,27 +135,18 @@ inline actor spawn_cl(const char* source,
* <tt>dims.empty()</tt>, or @p clCreateKernel
* <tt>dims.empty()</tt>, or @p clCreateKernel
* failed.
* failed.
*/
*/
template
<
typename
MapArgs
,
typename
MapResult
>
template
<
typename
MapArgs
,
typename
MapResult
>
inline
actor
spawn_cl
(
const
opencl
::
program
&
prog
,
inline
actor
const
char
*
fname
,
spawn_cl
(
const
opencl
::
program
&
prog
,
const
char
*
fname
,
MapArgs
map_args
,
MapArgs
map_args
,
MapResult
map_result
,
const
opencl
::
dim_vec
&
dims
,
MapResult
map_result
,
const
opencl
::
dim_vec
&
dims
,
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
local_dims
=
{},
const
opencl
::
dim_vec
&
local_dims
=
{},
size_t
result_size
=
0
)
{
size_t
result_size
=
0
)
{
using
std
::
move
;
using
std
::
move
;
typedef
typename
util
::
get_callable_trait
<
MapArgs
>::
fun_type
f0
;
typedef
typename
util
::
get_callable_trait
<
MapArgs
>::
fun_type
f0
;
typedef
typename
util
::
get_callable_trait
<
MapResult
>::
fun_type
f1
;
typedef
typename
util
::
get_callable_trait
<
MapResult
>::
fun_type
f1
;
detail
::
cl_spawn_helper
<
f0
,
f1
>
f
;
detail
::
cl_spawn_helper
<
f0
,
f1
>
f
;
return
f
(
f0
{
move
(
map_args
)},
return
f
(
f0
{
move
(
map_args
)},
f1
{
move
(
map_result
)},
prog
,
fname
,
dims
,
offset
,
f1
{
move
(
map_result
)},
local_dims
,
result_size
);
prog
,
fname
,
dims
,
offset
,
local_dims
,
result_size
);
}
}
/**
/**
...
@@ -182,24 +158,15 @@ inline actor spawn_cl(const opencl::program& prog,
...
@@ -182,24 +158,15 @@ inline actor spawn_cl(const opencl::program& prog,
* <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
<
typename
MapArgs
,
typename
MapResult
>
template
<
typename
MapArgs
,
typename
MapResult
>
inline
actor
spawn_cl
(
const
char
*
source
,
inline
actor
const
char
*
fun_name
,
spawn_cl
(
const
char
*
source
,
const
char
*
fun_name
,
MapArgs
map_args
,
MapArgs
map_args
,
MapResult
map_result
,
const
opencl
::
dim_vec
&
dims
,
MapResult
map_result
,
const
opencl
::
dim_vec
&
dims
,
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
offset
=
{},
const
opencl
::
dim_vec
&
local_dims
=
{},
const
opencl
::
dim_vec
&
local_dims
=
{},
size_t
result_size
=
0
)
{
size_t
result_size
=
0
)
{
using
std
::
move
;
using
std
::
move
;
return
spawn_cl
(
opencl
::
program
::
create
(
source
),
return
spawn_cl
(
opencl
::
program
::
create
(
source
),
fun_name
,
move
(
map_args
),
fun_name
,
move
(
map_result
),
dims
,
offset
,
local_dims
,
result_size
);
move
(
map_args
),
move
(
map_result
),
dims
,
offset
,
local_dims
,
result_size
);
}
}
}
// namespace caf
}
// namespace caf
...
...
libcaf_opencl/src/global.cpp
View file @
86e60176
...
@@ -118,13 +118,13 @@ std::string get_opencl_error(cl_int err) {
...
@@ -118,13 +118,13 @@ std::string get_opencl_error(cl_int err) {
return
"CL_INVALID_MIP_LEVEL"
;
return
"CL_INVALID_MIP_LEVEL"
;
case
CL_INVALID_GLOBAL_WORK_SIZE
:
case
CL_INVALID_GLOBAL_WORK_SIZE
:
return
"CL_INVALID_GLOBAL_WORK_SIZE"
;
return
"CL_INVALID_GLOBAL_WORK_SIZE"
;
default:
return
"UNKNOWN_ERROR"
;
default:
return
"UNKNOWN_ERROR"
;
}
}
}
}
cl_int
clReleaseDeviceDummy
(
cl_device_id
)
{
return
0
;
}
cl_int
clReleaseDeviceDummy
(
cl_device_id
)
{
return
0
;
}
cl_int
clRetainDeviceDummy
(
cl_device_id
)
{
return
0
;
}
cl_int
clRetainDeviceDummy
(
cl_device_id
)
{
return
0
;
}
}
// namespace opencl
}
// namespace opencl
}
// namespace caf
}
// namespace caf
libcaf_opencl/src/opencl_metainfo.cpp
View file @
86e60176
...
@@ -28,11 +28,9 @@ const std::vector<device_info> opencl_metainfo::get_devices() const {
...
@@ -28,11 +28,9 @@ const std::vector<device_info> opencl_metainfo::get_devices() const {
return
m_devices
;
return
m_devices
;
}
}
void
opencl_metainfo
::
initialize
()
void
opencl_metainfo
::
initialize
()
{
{
cl_int
err
{
0
};
cl_int
err
{
0
};
// get number of available platforms
// get number of available platforms
cl_uint
number_of_platforms
;
cl_uint
number_of_platforms
;
err
=
clGetPlatformIDs
(
0
,
nullptr
,
&
number_of_platforms
);
err
=
clGetPlatformIDs
(
0
,
nullptr
,
&
number_of_platforms
);
...
@@ -44,19 +42,16 @@ void opencl_metainfo::initialize()
...
@@ -44,19 +42,16 @@ void opencl_metainfo::initialize()
throw
logic_error
(
oss
.
str
());
throw
logic_error
(
oss
.
str
());
}
}
// get platform ids
// get platform ids
vector
<
cl_platform_id
>
ids
(
number_of_platforms
);
vector
<
cl_platform_id
>
ids
(
number_of_platforms
);
err
=
clGetPlatformIDs
(
ids
.
size
(),
ids
.
data
(),
nullptr
);
err
=
clGetPlatformIDs
(
ids
.
size
(),
ids
.
data
(),
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"clGetPlatformIDs (getting platform ids): "
oss
<<
"clGetPlatformIDs (getting platform ids): "
<<
get_opencl_error
(
err
);
<<
get_opencl_error
(
err
);
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
throw
logic_error
(
oss
.
str
());
throw
logic_error
(
oss
.
str
());
}
}
// find gpu devices on our platform
// find gpu devices on our platform
int
pid
{
0
};
int
pid
{
0
};
cl_uint
num_devices
{
0
};
cl_uint
num_devices
{
0
};
...
@@ -75,7 +70,8 @@ void opencl_metainfo::initialize()
...
@@ -75,7 +70,8 @@ void opencl_metainfo::initialize()
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
vector
<
cl_device_id
>
devices
(
num_devices
);
vector
<
cl_device_id
>
devices
(
num_devices
);
err
=
clGetDeviceIDs
(
ids
[
pid
],
dev_type
,
num_devices
,
devices
.
data
(),
nullptr
);
err
=
clGetDeviceIDs
(
ids
[
pid
],
dev_type
,
num_devices
,
devices
.
data
(),
nullptr
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"clGetDeviceIDs: "
<<
get_opencl_error
(
err
);
oss
<<
"clGetDeviceIDs: "
<<
get_opencl_error
(
err
);
...
@@ -83,24 +79,16 @@ void opencl_metainfo::initialize()
...
@@ -83,24 +79,16 @@ void opencl_metainfo::initialize()
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
auto
pfn_notify
=
[](
const
char
*
errinfo
,
auto
pfn_notify
=
[](
const
char
*
errinfo
,
const
void
*
,
size_t
,
void
*
)
{
const
void
*
,
CPPA_LOGC_ERROR
(
"cppa::opencl::opencl_metainfo"
,
"initialize"
,
size_t
,
void
*
)
{
CPPA_LOGC_ERROR
(
"cppa::opencl::opencl_metainfo"
,
"initialize"
,
"
\n
##### Error message via pfn_notify #####
\n
"
+
"
\n
##### Error message via pfn_notify #####
\n
"
+
string
(
errinfo
)
+
string
(
errinfo
)
+
"
\n
########################################"
);
"
\n
########################################"
);
};
};
// create a context
// create a context
m_context
.
adopt
(
clCreateContext
(
0
,
m_context
.
adopt
(
clCreateContext
(
0
,
devices
.
size
(),
devices
.
data
(),
pfn_notify
,
devices
.
size
(),
nullptr
,
&
err
));
devices
.
data
(),
pfn_notify
,
nullptr
,
&
err
));
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"clCreateContext: "
<<
get_opencl_error
(
err
);
oss
<<
"clCreateContext: "
<<
get_opencl_error
(
err
);
...
@@ -108,7 +96,6 @@ void opencl_metainfo::initialize()
...
@@ -108,7 +96,6 @@ void opencl_metainfo::initialize()
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
for
(
auto
&
d
:
devices
)
{
for
(
auto
&
d
:
devices
)
{
CPPA_LOG_TRACE
(
"Creating command queue for device(s)."
);
CPPA_LOG_TRACE
(
"Creating command queue for device(s)."
);
device_ptr
device
;
device_ptr
device
;
...
@@ -116,69 +103,54 @@ void opencl_metainfo::initialize()
...
@@ -116,69 +103,54 @@ void opencl_metainfo::initialize()
size_t
return_size
{
0
};
size_t
return_size
{
0
};
static
constexpr
size_t
buf_size
=
128
;
static
constexpr
size_t
buf_size
=
128
;
char
buf
[
buf_size
];
char
buf
[
buf_size
];
err
=
clGetDeviceInfo
(
device
.
get
(),
CL_DEVICE_NAME
,
buf_size
,
buf
,
&
return_size
);
err
=
clGetDeviceInfo
(
device
.
get
(),
CL_DEVICE_NAME
,
buf_size
,
buf
,
&
return_size
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_ERROR
,
"clGetDeviceInfo (CL_DEVICE_NAME): "
CPPA_LOGMF
(
CPPA_ERROR
,
<<
get_opencl_error
(
err
));
"clGetDeviceInfo (CL_DEVICE_NAME): "
<<
get_opencl_error
(
err
));
fill
(
buf
,
buf
+
buf_size
,
0
);
fill
(
buf
,
buf
+
buf_size
,
0
);
}
}
command_queue_ptr
cmd_queue
;
command_queue_ptr
cmd_queue
;
cmd_queue
.
adopt
(
clCreateCommandQueue
(
m_context
.
get
(),
cmd_queue
.
adopt
(
clCreateCommandQueue
(
m_context
.
get
(),
device
.
get
(),
device
.
get
(),
CL_QUEUE_PROFILING_ENABLE
,
&
err
));
CL_QUEUE_PROFILING_ENABLE
,
&
err
));
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
CPPA_LOGMF
(
CPPA_DEBUG
,
"Could not create command queue for device "
CPPA_LOGMF
(
CPPA_DEBUG
,
"Could not create command queue for device "
<<
buf
<<
": "
<<
get_opencl_error
(
err
));
<<
buf
<<
": "
<<
get_opencl_error
(
err
));
}
}
else
{
else
{
size_t
max_work_group_size
{
0
};
size_t
max_work_group_size
{
0
};
err
=
clGetDeviceInfo
(
device
.
get
(),
err
=
clGetDeviceInfo
(
device
.
get
(),
CL_DEVICE_MAX_WORK_GROUP_SIZE
,
CL_DEVICE_MAX_WORK_GROUP_SIZE
,
sizeof
(
size_t
),
&
max_work_group_size
,
&
return_size
);
sizeof
(
size_t
),
&
max_work_group_size
,
&
return_size
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"clGetDeviceInfo ("
oss
<<
"clGetDeviceInfo ("
<<
"CL_DEVICE_MAX_WORK_GROUP_SIZE): "
<<
"CL_DEVICE_MAX_WORK_GROUP_SIZE): "
<<
get_opencl_error
(
err
);
<<
get_opencl_error
(
err
);
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
cl_uint
max_work_item_dimensions
=
0
;
cl_uint
max_work_item_dimensions
=
0
;
err
=
clGetDeviceInfo
(
device
.
get
(),
err
=
clGetDeviceInfo
(
device
.
get
(),
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
,
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS
,
sizeof
(
cl_uint
),
&
max_work_item_dimensions
,
sizeof
(
cl_uint
),
&
max_work_item_dimensions
,
&
return_size
);
&
return_size
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"clGetDeviceInfo ("
oss
<<
"clGetDeviceInfo ("
<<
"CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): "
<<
"CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS): "
<<
get_opencl_error
(
err
);
<<
get_opencl_error
(
err
);
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
dim_vec
max_work_items_per_dim
(
max_work_item_dimensions
);
dim_vec
max_work_items_per_dim
(
max_work_item_dimensions
);
err
=
clGetDeviceInfo
(
device
.
get
(),
err
=
clGetDeviceInfo
(
device
.
get
(),
CL_DEVICE_MAX_WORK_ITEM_SIZES
,
CL_DEVICE_MAX_WORK_ITEM_SIZES
,
sizeof
(
size_t
)
*
max_work_item_dimensions
,
sizeof
(
size_t
)
*
max_work_item_dimensions
,
max_work_items_per_dim
.
data
(),
&
return_size
);
max_work_items_per_dim
.
data
(),
&
return_size
);
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"clGetDeviceInfo ("
oss
<<
"clGetDeviceInfo ("
<<
"CL_DEVICE_MAX_WORK_ITEM_SIZES): "
<<
"CL_DEVICE_MAX_WORK_ITEM_SIZES): "
<<
get_opencl_error
(
err
);
<<
get_opencl_error
(
err
);
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
CPPA_LOGMF
(
CPPA_ERROR
,
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
device_info
dev_info
{
device
,
device_info
dev_info
{
device
,
cmd_queue
,
max_work_group_size
,
cmd_queue
,
max_work_item_dimensions
,
max_work_items_per_dim
};
max_work_group_size
,
max_work_item_dimensions
,
max_work_items_per_dim
};
m_devices
.
push_back
(
move
(
dev_info
));
m_devices
.
push_back
(
move
(
dev_info
));
}
}
}
}
...
@@ -192,13 +164,9 @@ void opencl_metainfo::initialize()
...
@@ -192,13 +164,9 @@ void opencl_metainfo::initialize()
}
}
}
}
void
opencl_metainfo
::
destroy
()
{
void
opencl_metainfo
::
destroy
()
{
delete
this
;
}
delete
this
;
}
void
opencl_metainfo
::
dispose
()
{
void
opencl_metainfo
::
dispose
()
{
delete
this
;
}
delete
this
;
}
opencl_metainfo
*
get_opencl_metainfo
()
{
opencl_metainfo
*
get_opencl_metainfo
()
{
return
detail
::
singleton_manager
::
get_opencl_metainfo
();
return
detail
::
singleton_manager
::
get_opencl_metainfo
();
...
@@ -206,5 +174,3 @@ opencl_metainfo* get_opencl_metainfo() {
...
@@ -206,5 +174,3 @@ opencl_metainfo* get_opencl_metainfo() {
}
// namespace opencl
}
// namespace opencl
}
// namespace caf
}
// namespace caf
libcaf_opencl/src/program.cpp
View file @
86e60176
...
@@ -31,21 +31,23 @@ using namespace std;
...
@@ -31,21 +31,23 @@ using namespace std;
namespace
caf
{
namespace
caf
{
namespace
opencl
{
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
)),
m_program
(
move
(
program
)),
m_queue
(
move
(
queue
))
{
}
:
m_context
(
move
(
context
))
,
m_program
(
move
(
program
))
program
program
::
create
(
const
char
*
kernel_source
,
const
char
*
options
,
uint32_t
device_id
)
{
,
m_queue
(
move
(
queue
))
{}
program
program
::
create
(
const
char
*
kernel_source
,
const
char
*
options
,
uint32_t
device_id
)
{
auto
metainfo
=
get_opencl_metainfo
();
auto
metainfo
=
get_opencl_metainfo
();
auto
devices
=
metainfo
->
get_devices
();
auto
devices
=
metainfo
->
get_devices
();
auto
context
=
metainfo
->
m_context
;
auto
context
=
metainfo
->
m_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
<<
" is not a vaild device. Maximum id is: "
<<
" is not a vaild device. Maximum id is: "
<<
(
devices
.
size
()
-
1
)
<<
(
devices
.
size
()
-
1
)
<<
"."
;
<<
"."
;
CPPA_LOGM_ERROR
(
detail
::
demangle
<
program
>
().
c_str
(),
oss
.
str
());
CPPA_LOGM_ERROR
(
detail
::
demangle
<
program
>
().
c_str
(),
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
throw
runtime_error
(
oss
.
str
());
}
}
...
@@ -55,14 +57,10 @@ program program::create(const char* kernel_source, const char* options, uint32_t
...
@@ -55,14 +57,10 @@ program program::create(const char* kernel_source, const char* options, uint32_t
// create program object from kernel source
// create program object from kernel source
size_t
kernel_source_length
=
strlen
(
kernel_source
);
size_t
kernel_source_length
=
strlen
(
kernel_source
);
program_ptr
pptr
;
program_ptr
pptr
;
pptr
.
adopt
(
clCreateProgramWithSource
(
context
.
get
(),
pptr
.
adopt
(
clCreateProgramWithSource
(
context
.
get
(),
1
,
&
kernel_source
,
1
,
&
kernel_source_length
,
&
err
));
&
kernel_source
,
&
kernel_source_length
,
&
err
));
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
throw
runtime_error
(
"clCreateProgramWithSource: "
throw
runtime_error
(
"clCreateProgramWithSource: "
+
get_opencl_error
(
err
));
+
get_opencl_error
(
err
));
}
}
// build programm from program object
// build programm from program object
...
@@ -71,34 +69,27 @@ program program::create(const char* kernel_source, const char* options, uint32_t
...
@@ -71,34 +69,27 @@ program program::create(const char* kernel_source, const char* options, uint32_t
if
(
err
!=
CL_SUCCESS
)
{
if
(
err
!=
CL_SUCCESS
)
{
ostringstream
oss
;
ostringstream
oss
;
oss
<<
"clBuildProgram: "
<<
get_opencl_error
(
err
);
oss
<<
"clBuildProgram: "
<<
get_opencl_error
(
err
);
// the build log will be printed by the
// the build log will be printed by the
// pfn_notify (see opencl_metainfo.cpp)
// pfn_notify (see opencl_metainfo.cpp)
#ifndef __APPLE__
#ifndef __APPLE__
// seems that just apple implemented the
// seems that just apple implemented the
// pfn_notify callback, but we can get
// pfn_notify callback, but we can get
// the build log
// the build log
if
(
err
==
CL_BUILD_PROGRAM_FAILURE
)
{
if
(
err
==
CL_BUILD_PROGRAM_FAILURE
)
{
size_t
buildlog_buffer_size
=
0
;
size_t
buildlog_buffer_size
=
0
;
// get the log length
// get the log length
clGetProgramBuildInfo
(
pptr
.
get
(),
clGetProgramBuildInfo
(
pptr
.
get
(),
dev_tmp
,
CL_PROGRAM_BUILD_LOG
,
dev_tmp
,
sizeof
(
buildlog_buffer_size
),
nullptr
,
CL_PROGRAM_BUILD_LOG
,
sizeof
(
buildlog_buffer_size
),
nullptr
,
&
buildlog_buffer_size
);
&
buildlog_buffer_size
);
vector
<
char
>
buffer
(
buildlog_buffer_size
);
vector
<
char
>
buffer
(
buildlog_buffer_size
);
// fill the buffer with buildlog informations
// fill the buffer with buildlog informations
clGetProgramBuildInfo
(
pptr
.
get
(),
clGetProgramBuildInfo
(
pptr
.
get
(),
dev_tmp
,
CL_PROGRAM_BUILD_LOG
,
dev_tmp
,
CL_PROGRAM_BUILD_LOG
,
sizeof
(
buffer
[
0
])
*
buildlog_buffer_size
,
sizeof
(
buffer
[
0
])
*
buildlog_buffer_size
,
buffer
.
data
(),
buffer
.
data
(),
nullptr
);
nullptr
);
CPPA_LOGC_ERROR
(
"cppa::opencl::program"
,
CPPA_LOGC_ERROR
(
"cppa::opencl::program"
,
"create"
,
"create"
,
"Build log:
\n
"
+
string
(
buffer
.
data
())
+
"Build log:
\n
"
+
string
(
buffer
.
data
())
+
"
\n
########################################"
);
"
\n
########################################"
);
}
}
...
@@ -110,4 +101,3 @@ program program::create(const char* kernel_source, const char* options, uint32_t
...
@@ -110,4 +101,3 @@ program program::create(const char* kernel_source, const char* options, uint32_t
}
// namespace opencl
}
// namespace opencl
}
// namespace caf
}
// namespace caf
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