From ad33665b8912a561fca2c9a847fe729f97cce446 Mon Sep 17 00:00:00 2001
From: Kyle Lutz <kyle.r.lutz@gmail.com>
Date: Tue, 30 Sep 2014 22:22:44 -0700
Subject: [PATCH] Add type_definition<T>() function

---
 include/boost/compute/type_traits.hpp         |  1 +
 .../compute/type_traits/type_definition.hpp   | 26 ++++++++++++
 include/boost/compute/types/struct.hpp        | 27 +++++++-----
 test/test_struct.cpp                          | 42 +++++++++++++++++++
 4 files changed, 86 insertions(+), 10 deletions(-)
 create mode 100644 include/boost/compute/type_traits/type_definition.hpp

diff --git a/include/boost/compute/type_traits.hpp b/include/boost/compute/type_traits.hpp
index 74c9fc7d..1673266d 100644
--- a/include/boost/compute/type_traits.hpp
+++ b/include/boost/compute/type_traits.hpp
@@ -16,6 +16,7 @@
 #include <boost/compute/type_traits/is_vector_type.hpp>
 #include <boost/compute/type_traits/make_vector_type.hpp>
 #include <boost/compute/type_traits/scalar_type.hpp>
+#include <boost/compute/type_traits/type_definition.hpp>
 #include <boost/compute/type_traits/type_name.hpp>
 #include <boost/compute/type_traits/vector_size.hpp>
 
diff --git a/include/boost/compute/type_traits/type_definition.hpp b/include/boost/compute/type_traits/type_definition.hpp
new file mode 100644
index 00000000..3e052f67
--- /dev/null
+++ b/include/boost/compute/type_traits/type_definition.hpp
@@ -0,0 +1,26 @@
+//---------------------------------------------------------------------------//
+// 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
+// http://www.boost.org/LICENSE_1_0.txt
+//
+// See http://kylelutz.github.com/compute for more information.
+//---------------------------------------------------------------------------//
+
+#ifndef BOOST_COMPUTE_TYPE_TRAITS_TYPE_DEFINITION_HPP
+#define BOOST_COMPUTE_TYPE_TRAITS_TYPE_DEFINITION_HPP
+
+#include <string>
+
+namespace boost {
+namespace compute {
+
+/// Returns the OpenCL type definition for \p T.
+template<class T>
+std::string type_definition();
+
+} // end compute namespace
+} // end boost namespace
+
+#endif // BOOST_COMPUTE_TYPE_TRAITS_TYPE_DEFINITION_HPP
diff --git a/include/boost/compute/types/struct.hpp b/include/boost/compute/types/struct.hpp
index 600c1904..12fa4ef3 100644
--- a/include/boost/compute/types/struct.hpp
+++ b/include/boost/compute/types/struct.hpp
@@ -20,6 +20,7 @@
 #include <boost/preprocessor/seq/fold_left.hpp>
 #include <boost/preprocessor/seq/transform.hpp>
 
+#include <boost/compute/type_traits/type_definition.hpp>
 #include <boost/compute/type_traits/type_name.hpp>
 #include <boost/compute/detail/meta_kernel.hpp>
 #include <boost/compute/detail/variadic_macros.hpp>
@@ -124,21 +125,27 @@ inline std::string adapt_struct_insert_member(T Struct::*, const char *name)
         "BOOST_COMPUTE_ADAPT_STRUCT() does not support structs with internal padding." \
     ); \
     BOOST_COMPUTE_TYPE_NAME(type, name) \
-    namespace boost { namespace compute { namespace detail { \
+    namespace boost { namespace compute { \
+    template<> \
+    inline std::string type_definition<type>() \
+    { \
+        std::stringstream declaration; \
+        declaration << "typedef struct __attribute__((packed)) {\n" \
+                    BOOST_PP_SEQ_FOR_EACH( \
+                        BOOST_COMPUTE_DETAIL_ADAPT_STRUCT_INSERT_MEMBER, \
+                        type, \
+                        BOOST_COMPUTE_PP_TUPLE_TO_SEQ(members) \
+                    ) \
+                    << "} " << type_name<type>() << ";\n"; \
+        return declaration.str(); \
+    } \
+    namespace detail { \
     template<> \
     struct inject_type_impl<type> \
     { \
         void operator()(meta_kernel &kernel) \
         { \
-            std::stringstream declaration; \
-            declaration << "typedef struct __attribute__((packed)) {\n" \
-                        BOOST_PP_SEQ_FOR_EACH( \
-                            BOOST_COMPUTE_DETAIL_ADAPT_STRUCT_INSERT_MEMBER, \
-                            type, \
-                            BOOST_COMPUTE_PP_TUPLE_TO_SEQ(members) \
-                        ) \
-                        << "} " << type_name<type>() << ";\n"; \
-            kernel.add_type_declaration<type>(declaration.str()); \
+            kernel.add_type_declaration<type>(type_definition<type>()); \
         } \
     }; \
     inline meta_kernel& operator<<(meta_kernel &k, type s) \
diff --git a/test/test_struct.cpp b/test/test_struct.cpp
index b4ed187c..2ccea288 100644
--- a/test/test_struct.cpp
+++ b/test/test_struct.cpp
@@ -14,11 +14,14 @@
 #include <boost/test/unit_test.hpp>
 
 #include <boost/compute/function.hpp>
+#include <boost/compute/source.hpp>
 #include <boost/compute/algorithm/find_if.hpp>
 #include <boost/compute/algorithm/transform.hpp>
 #include <boost/compute/container/vector.hpp>
 #include <boost/compute/functional/field.hpp>
 #include <boost/compute/types/struct.hpp>
+#include <boost/compute/type_traits/type_name.hpp>
+#include <boost/compute/type_traits/type_definition.hpp>
 
 namespace compute = boost::compute;
 
@@ -82,4 +85,43 @@ BOOST_AUTO_TEST_CASE(atom_struct)
     CHECK_RANGE_EQUAL(int, 3, atomic_numbers, (1, 1, 8));
 }
 
+BOOST_AUTO_TEST_CASE(custom_kernel)
+{
+    std::vector<chemistry::Atom> data;
+    data.push_back(chemistry::Atom(1.f, 0.f, 0.f, 1));
+    data.push_back(chemistry::Atom(0.f, 1.f, 0.f, 1));
+    data.push_back(chemistry::Atom(0.f, 0.f, 0.f, 8));
+
+    compute::vector<chemistry::Atom> atoms(data.size(), context);
+    compute::copy(data.begin(), data.end(), atoms.begin(), queue);
+
+    std::string source = BOOST_COMPUTE_STRINGIZE_SOURCE(
+        __kernel void custom_kernel(__global const Atom *atoms,
+                                    __global float *distances)
+        {
+            const uint i = get_global_id(0);
+            const __global Atom *atom = &atoms[i];
+
+            const float4 center = { 0, 0, 0, 0 };
+            const float4 position = { atom->x, atom->y, atom->z, 0 };
+
+            distances[i] = distance(position, center);
+        }
+    );
+
+    // add type definition for Atom to the start of the program source
+    source = compute::type_definition<chemistry::Atom>() + "\n" + source;
+
+    compute::program program =
+        compute::program::build_with_source(source, context);
+
+    compute::vector<float> distances(atoms.size(), context);
+
+    compute::kernel custom_kernel = program.create_kernel("custom_kernel");
+    custom_kernel.set_arg(0, atoms);
+    custom_kernel.set_arg(1, distances);
+
+    queue.enqueue_1d_range_kernel(custom_kernel, 0, atoms.size(), 1);
+}
+
 BOOST_AUTO_TEST_SUITE_END()