diff --git a/include/boost/compute/detail/get_object_info.hpp b/include/boost/compute/detail/get_object_info.hpp
index dca7f1e8..1c54ed49 100644
--- a/include/boost/compute/detail/get_object_info.hpp
+++ b/include/boost/compute/detail/get_object_info.hpp
@@ -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.
diff --git a/include/boost/compute/kernel.hpp b/include/boost/compute/kernel.hpp
index 49c7a166..e588f852 100644
--- a/include/boost/compute/kernel.hpp
+++ b/include/boost/compute/kernel.hpp
@@ -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.
diff --git a/include/boost/compute/program.hpp b/include/boost/compute/program.hpp
index 56c42955..81949ae5 100644
--- a/include/boost/compute/program.hpp
+++ b/include/boost/compute/program.hpp
@@ -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.
diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt
index 9a86f6ba..195acb72 100644
--- a/test/CMakeLists.txt
+++ b/test/CMakeLists.txt
@@ -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)
diff --git a/test/test_amd_cpp_kernel_language.cpp b/test/test_amd_cpp_kernel_language.cpp
index 0a3d82df..bd739315 100644
--- a/test/test_amd_cpp_kernel_language.cpp
+++ b/test/test_amd_cpp_kernel_language.cpp
@@ -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>
diff --git a/test/test_command_queue.cpp b/test/test_command_queue.cpp
index 4a95fb88..028e9014 100644
--- a/test/test_command_queue.cpp
+++ b/test/test_command_queue.cpp
@@ -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>
diff --git a/test/test_generate.cpp b/test/test_generate.cpp
index 95d617a2..7212efc5 100644
--- a/test/test_generate.cpp
+++ b/test/test_generate.cpp
@@ -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>
diff --git a/test/test_image2d.cpp b/test/test_image2d.cpp
index fb96c1b8..fc4e5a5e 100644
--- a/test/test_image2d.cpp
+++ b/test/test_image2d.cpp
@@ -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>
diff --git a/test/test_image_sampler.cpp b/test/test_image_sampler.cpp
index 452859eb..6fe74ec4 100644
--- a/test/test_image_sampler.cpp
+++ b/test/test_image_sampler.cpp
@@ -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>
 
diff --git a/test/test_kernel.cpp b/test/test_kernel.cpp
index ce2b0aa1..0c6c6126 100644
--- a/test/test_kernel.cpp
+++ b/test/test_kernel.cpp
@@ -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()
diff --git a/test/test_pair.cpp b/test/test_pair.cpp
index 82ca3884..3eef7732 100644
--- a/test/test_pair.cpp
+++ b/test/test_pair.cpp
@@ -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>
diff --git a/test/test_program.cpp b/test/test_program.cpp
index c3895fe6..83b44d3a 100644
--- a/test/test_program.cpp
+++ b/test/test_program.cpp
@@ -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()
diff --git a/test/test_tuple.cpp b/test/test_tuple.cpp
index 4c36b6a8..04fec7cd 100644
--- a/test/test_tuple.cpp
+++ b/test/test_tuple.cpp
@@ -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>