boost/compute
2
0
Fork 0
mirror of https://github.com/boostorg/compute.git synced 2026-01-27 18:52:15 +00:00
Code Issues Projects Releases Wiki Activity

Merge pull request #264 from kylelutz/refactor-get-info

Browse Source 
Refactor get_info() functions
This commit is contained in:
Kyle Lutz
2014-09-14 21:13:36 -07:00
parent b95763d8b8 cdcd4c5a32
commit 4b10ea608b
13 changed files with 209 additions and 141 deletions
Download Patch File Download Diff File Expand all files Collapse all files

138
include/boost/compute/detail/get_object_info.hpp
View File

@@ -1,5 +1,5 @@
//---------------------------------------------------------------------------//
// Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
// Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@gmail.com>
//
// Distributed under the Boost Software License, Version 1.0
// See accompanying file LICENSE_1_0.txt or copy at
@@ -16,24 +16,77 @@
#include <boost/throw_exception.hpp>
#include <boost/compute/exception.hpp>
#include <boost/compute/cl.hpp>
#include <boost/compute/exception/opencl_error.hpp>
namespace boost {
namespace compute {
namespace detail {
// default implementation
template<class T, class Function, class Object, class Info>
struct _get_object_info_impl
template<class Function, class Object, class AuxInfo>
struct bound_info_function
{
T operator()(Function function, Object object, Info info)
bound_info_function(Function function, Object object, AuxInfo aux_info)
: m_function(function),
m_object(object),
m_aux_info(aux_info)
{
}
template<class Info>
cl_int operator()(Info info, size_t size, void *value, size_t *size_ret) const
{
return m_function(m_object, m_aux_info, info, size, value, size_ret);
}
Function m_function;
Object m_object;
AuxInfo m_aux_info;
};
template<class Function, class Object>
struct bound_info_function<Function, Object, void>
{
bound_info_function(Function function, Object object)
: m_function(function),
m_object(object)
{
}
template<class Info>
cl_int operator()(Info info, size_t size, void *value, size_t *size_ret) const
{
return m_function(m_object, info, size, value, size_ret);
}
Function m_function;
Object m_object;
};
template<class Function, class Object>
inline bound_info_function<Function, Object, void>
bind_info_function(Function f, Object o)
{
return bound_info_function<Function, Object, void>(f, o);
}
template<class Function, class Object, class AuxInfo>
inline bound_info_function<Function, Object, AuxInfo>
bind_info_function(Function f, Object o, AuxInfo j)
{
return bound_info_function<Function, Object, AuxInfo>(f, o, j);
}
// default implementation
template<class T>
struct get_object_info_impl
{
template<class Function, class Info>
T operator()(Function function, Info info) const
{
T value;
cl_int ret = function(object,
static_cast<cl_uint>(info),
sizeof(T),
&value,
0);
cl_int ret = function(info, sizeof(T), &value, 0);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
@@ -43,17 +96,15 @@ struct _get_object_info_impl
};
// specialization for bool
template<class Function, class Object, class Info>
struct _get_object_info_impl<bool, Function, Object, Info>
template<>
struct get_object_info_impl<bool>
{
bool operator()(Function function, Object object, Info info)
template<class Function, class Info>
bool operator()(Function function, Info info) const
{
cl_bool value;
cl_int ret = function(object,
static_cast<cl_uint>(info),
sizeof(cl_bool),
&value,
0);
cl_int ret = function(info, sizeof(cl_bool), &value, 0);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
@@ -63,17 +114,15 @@ struct _get_object_info_impl<bool, Function, Object, Info>
};
// specialization for std::string
template<class Function, class Object, class Info>
struct _get_object_info_impl<std::string, Function, Object, Info>
template<>
struct get_object_info_impl<std::string>
{
std::string operator()(Function function, Object object, Info info)
template<class Function, class Info>
std::string operator()(Function function, Info info) const
{
size_t size = 0;
cl_int ret = function(object,
static_cast<cl_uint>(info),
0,
0,
&size);
cl_int ret = function(info, 0, 0, &size);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
@@ -83,11 +132,8 @@ struct _get_object_info_impl<std::string, Function, Object, Info>
}
std::string value(size - 1, 0);
ret = function(object,
static_cast<cl_uint>(info),
size,
&value[0],
0);
ret = function(info, size, &value[0], 0);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
@@ -97,27 +143,21 @@ struct _get_object_info_impl<std::string, Function, Object, Info>
};
// specialization for std::vector<T>
template<class T, class Function, class Object, class Info>
struct _get_object_info_impl<std::vector<T>, Function, Object, Info>
template<class T>
struct get_object_info_impl<std::vector<T> >
{
std::vector<T> operator()(Function function, Object object, Info info)
template<class Function, class Info>
std::vector<T> operator()(Function function, Info info) const
{
size_t size = 0;
cl_int ret = function(object,
static_cast<cl_uint>(info),
0,
0,
&size);
cl_int ret = function(info, 0, 0, &size);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
std::vector<T> vector(size / sizeof(T));
ret = function(object,
static_cast<cl_uint>(info),
size,
&vector[0],
0);
ret = function(info, size, &vector[0], 0);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
@@ -130,7 +170,13 @@ struct _get_object_info_impl<std::vector<T>, Function, Object, Info>
template<class T, class Function, class Object, class Info>
inline T get_object_info(Function f, Object o, Info i)
{
return _get_object_info_impl<T, Function, Object, Info>()(f, o, i);
return get_object_info_impl<T>()(bind_info_function(f, o), i);
}
template<class T, class Function, class Object, class Info, class AuxInfo>
inline T get_object_info(Function f, Object o, Info i, AuxInfo j)
{
return get_object_info_impl<T>()(bind_info_function(f, o, j), i);
}
// returns the value type for the clGet*Info() call on Object with Enum.

28
include/boost/compute/kernel.hpp
View File

@@ -175,20 +175,9 @@ public:
///
/// \see_opencl_ref{clGetKernelArgInfo}
template<class T>
T get_arg_info(size_t index, cl_kernel_arg_info info)
T get_arg_info(size_t index, cl_kernel_arg_info info) const
{
T value;
cl_int ret = clGetKernelArgInfo(m_kernel,
index,
info,
sizeof(T),
&value,
0);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return value;
return detail::get_object_info<T>(clGetKernelArgInfo, m_kernel, info, index);
}
#endif // CL_VERSION_1_2
@@ -198,18 +187,7 @@ public:
template<class T>
T get_work_group_info(const device &device, cl_kernel_work_group_info info)
{
T value;
cl_int ret = clGetKernelWorkGroupInfo(m_kernel,
device.id(),
info,
sizeof(T),
&value,
0);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return value;
return detail::get_object_info<T>(clGetKernelWorkGroupInfo, m_kernel, info, device.id());
}
/// Sets the argument at \p index to \p value with \p size.

57
include/boost/compute/program.hpp
View File

@@ -181,19 +181,8 @@ public:
std::vector<device> get_devices() const
{
size_t device_ids_size = 0;
clGetProgramInfo(m_program,
CL_PROGRAM_DEVICES,
0,
0,
&device_ids_size);
std::vector<cl_device_id> device_ids(device_ids_size / sizeof(size_t));
clGetProgramInfo(m_program,
CL_PROGRAM_DEVICES,
device_ids_size,
&device_ids[0],
0);
std::vector<cl_device_id> device_ids =
get_info<std::vector<cl_device_id> >(CL_PROGRAM_DEVICES);
std::vector<device> devices;
for(size_t i = 0; i < device_ids.size(); i++){
@@ -223,6 +212,22 @@ public:
typename detail::get_object_info_type<program, Enum>::type
get_info() const;
/// Returns build information about the program.
///
/// For example, this function can be used to retreive the options used
/// to build the program:
/// \code
/// std::string build_options =
/// program.get_build_info<std::string>(CL_PROGRAM_BUILD_OPTIONS);
/// \endcode
///
/// \see_opencl_ref{clGetProgramInfo}
template<class T>
T get_build_info(cl_program_build_info info, const device &device) const
{
return detail::get_object_info<T>(clGetProgramBuildInfo, m_program, info, device.id());
}
/// Builds the program with \p options.
///
/// If the program fails to compile, this function will throw an
@@ -329,31 +334,7 @@ public:
/// Returns the build log.
std::string build_log() const
{
device device = get_devices()[0];
size_t size = 0;
cl_int ret = clGetProgramBuildInfo(m_program,
device.id(),
CL_PROGRAM_BUILD_LOG,
0,
0,
&size);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
std::string value(size - 1, 0);
ret = clGetProgramBuildInfo(m_program,
device.id(),
CL_PROGRAM_BUILD_LOG,
size,
&value[0],
0);
if(ret != CL_SUCCESS){
BOOST_THROW_EXCEPTION(opencl_error(ret));
}
return value;
return get_build_info<std::string>(CL_PROGRAM_BUILD_LOG, get_devices().front());
}
/// Creates and returns a new kernel object for \p name.

2
test/CMakeLists.txt
View File

@@ -10,8 +10,6 @@ endif()
find_package(Boost 1.48 REQUIRED COMPONENTS ${BOOST_COMPONENTS})
add_definitions(-DBOOST_TEST_DYN_LINK)
add_definitions(-DBOOST_COMPUTE_DEBUG_KERNEL_COMPILATION)
# enable code coverage generation (only with GCC for now)
if(${BOOST_COMPUTE_ENABLE_COVERAGE} AND "${CMAKE_CXX_COMPILER_ID}" STREQUAL "GNU")
add_definitions(-fprofile-arcs -ftest-coverage)

2
test/test_amd_cpp_kernel_language.cpp
View File

@@ -11,6 +11,8 @@
#define BOOST_TEST_MODULE TestAmdCppKernel
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/compute/kernel.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/source.hpp>

2
test/test_command_queue.cpp
View File

@@ -11,6 +11,8 @@
#define BOOST_TEST_MODULE TestCommandQueue
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/compute/kernel.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/program.hpp>

2
test/test_generate.cpp
View File

@@ -11,6 +11,8 @@
#define BOOST_TEST_MODULE TestGenerate
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/compute/function.hpp>
#include <boost/compute/algorithm/generate.hpp>
#include <boost/compute/algorithm/generate_n.hpp>

2
test/test_image2d.cpp
View File

@@ -11,6 +11,8 @@
#define BOOST_TEST_MODULE TestImage2D
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/compute/system.hpp>
#include <boost/compute/image2d.hpp>
#include <boost/compute/algorithm/find.hpp>

2
test/test_image_sampler.cpp
View File

@@ -11,6 +11,8 @@
#define BOOST_TEST_MODULE TestImageSampler
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/compute/system.hpp>
#include <boost/compute/image_sampler.hpp>

91
test/test_kernel.cpp
View File

@@ -18,36 +18,36 @@
#include "context_setup.hpp"
namespace bc = boost::compute;
namespace compute = boost::compute;
BOOST_AUTO_TEST_CASE(name)
{
bc::kernel foo = bc::kernel::create_with_source("__kernel void foo(int x) { }",
"foo",
context);
compute::kernel foo = compute::kernel::create_with_source(
"__kernel void foo(int x) { }", "foo", context
);
BOOST_CHECK_EQUAL(foo.name(), "foo");
bc::kernel bar = bc::kernel::create_with_source("__kernel void bar(float x) { }",
"bar",
context);
compute::kernel bar = compute::kernel::create_with_source(
"__kernel void bar(float x) { }", "bar", context
);
BOOST_CHECK_EQUAL(bar.name(), "bar");
}
BOOST_AUTO_TEST_CASE(arity)
{
bc::kernel foo = bc::kernel::create_with_source("__kernel void foo(int x) { }",
"foo",
context);
compute::kernel foo = compute::kernel::create_with_source(
"__kernel void foo(int x) { }", "foo", context
);
BOOST_CHECK_EQUAL(foo.arity(), size_t(1));
bc::kernel bar = bc::kernel::create_with_source("__kernel void bar(float x, float y) { }",
"bar",
context);
compute::kernel bar = compute::kernel::create_with_source(
"__kernel void bar(float x, float y) { }", "bar", context
);
BOOST_CHECK_EQUAL(bar.arity(), size_t(2));
bc::kernel baz = bc::kernel::create_with_source("__kernel void baz(char x, char y, char z) { }",
"baz",
context);
compute::kernel baz = compute::kernel::create_with_source(
"__kernel void baz(char x, char y, char z) { }", "baz", context
);
BOOST_CHECK_EQUAL(baz.arity(), size_t(3));
}
@@ -60,11 +60,11 @@ BOOST_AUTO_TEST_CASE(set_buffer_arg)
}
);
bc::kernel foo =
bc::kernel::create_with_source(source, "foo", context);
compute::kernel foo =
compute::kernel::create_with_source(source, "foo", context);
bc::buffer x(context, 16);
bc::buffer y(context, 16);
compute::buffer x(context, 16);
compute::buffer y(context, 16);
foo.set_arg(0, x);
foo.set_arg(1, y.get());
@@ -83,14 +83,14 @@ BOOST_AUTO_TEST_CASE(get_work_group_info)
" scratch[lid] = input[gid];\n"
"}\n";
boost::compute::program program =
boost::compute::program::create_with_source(source, context);
compute::program program =
compute::program::create_with_source(source, context);
program.build();
boost::compute::kernel kernel = program.create_kernel("sum");
compute::kernel kernel = program.create_kernel("sum");
using boost::compute::ulong_;
using compute::ulong_;
// get local memory size
kernel.get_work_group_info<ulong_>(device, CL_KERNEL_LOCAL_MEM_SIZE);
@@ -104,15 +104,46 @@ BOOST_AUTO_TEST_CASE(get_work_group_info)
#ifndef BOOST_NO_VARIADIC_TEMPLATES
BOOST_AUTO_TEST_CASE(kernel_set_args)
{
bc::kernel k =
bc::kernel::create_with_source(
"__kernel void test(int x, float y, char z) { }",
"test",
context
);
compute::kernel k = compute::kernel::create_with_source(
"__kernel void test(int x, float y, char z) { }", "test", context
);
k.set_args(4, 2.4f, 'a');
}
#endif
#ifdef CL_VERSION_1_2
BOOST_AUTO_TEST_CASE(get_arg_info)
{
const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
__kernel void sum_kernel(__global const int *input,
const uint size,
__global int *result)
{
int sum = 0;
for(uint i = 0; i < size; i++){
sum += input[i];
}
*result = sum;
}
);
compute::program program =
compute::program::create_with_source(source, context);
program.build("-cl-kernel-arg-info");
compute::kernel kernel = program.create_kernel("sum_kernel");
BOOST_CHECK_EQUAL(kernel.get_info<CL_KERNEL_NUM_ARGS>(), 3);
BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_TYPE_NAME), "int*");
BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_NAME), "input");
BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_TYPE_NAME), "uint");
BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_NAME), "size");
BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_TYPE_NAME), "int*");
BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_NAME), "result");
}
#endif // CL_VERSION_1_2
BOOST_AUTO_TEST_SUITE_END()

2
test/test_pair.cpp
View File

@@ -11,6 +11,8 @@
#define BOOST_TEST_MODULE TestPair
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/compute/algorithm/copy.hpp>
#include <boost/compute/algorithm/fill.hpp>
#include <boost/compute/algorithm/find.hpp>

20
test/test_program.cpp
View File

@@ -155,4 +155,24 @@ BOOST_AUTO_TEST_CASE(compile_and_link)
}
#endif // CL_VERSION_1_2
BOOST_AUTO_TEST_CASE(build_log)
{
const char invalid_source[] =
"__kernel void foo(__global int *input) { !@#$%^&*() }";
compute::program invalid_program =
compute::program::create_with_source(invalid_source, context);
try {
invalid_program.build();
// should not get here
BOOST_CHECK(false);
}
catch(compute::opencl_error &e){
std::string log = invalid_program.build_log();
BOOST_CHECK(!log.empty());
}
}
BOOST_AUTO_TEST_SUITE_END()

2
test/test_tuple.cpp
View File

@@ -11,6 +11,8 @@
#define BOOST_TEST_MODULE TestTuple
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/tuple/tuple.hpp>
#include <boost/tuple/tuple_io.hpp>
#include <boost/tuple/tuple_comparison.hpp>