1 //---------------------------------------------------------------------------//
2 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
4 // Distributed under the Boost Software License, Version 1.0
5 // See accompanying file LICENSE_1_0.txt or copy at
6 // http://www.boost.org/LICENSE_1_0.txt
8 // See http://boostorg.github.com/compute for more information.
9 //---------------------------------------------------------------------------//
11 #ifndef BOOST_COMPUTE_DETAIL_META_KERNEL_HPP
12 #define BOOST_COMPUTE_DETAIL_META_KERNEL_HPP
21 #include <boost/tuple/tuple.hpp>
22 #include <boost/type_traits.hpp>
23 #include <boost/lexical_cast.hpp>
24 #include <boost/static_assert.hpp>
25 #include <boost/algorithm/string/find.hpp>
26 #include <boost/preprocessor/repetition.hpp>
28 #include <boost/compute/kernel.hpp>
29 #include <boost/compute/closure.hpp>
30 #include <boost/compute/function.hpp>
31 #include <boost/compute/functional.hpp>
32 #include <boost/compute/type_traits.hpp>
33 #include <boost/compute/command_queue.hpp>
34 #include <boost/compute/image/image2d.hpp>
35 #include <boost/compute/image/image_sampler.hpp>
36 #include <boost/compute/memory_object.hpp>
37 #include <boost/compute/memory/svm_ptr.hpp>
38 #include <boost/compute/detail/device_ptr.hpp>
39 #include <boost/compute/detail/sha1.hpp>
40 #include <boost/compute/utility/program_cache.hpp>
47 class meta_kernel_variable
50 typedef T result_type;
52 meta_kernel_variable(const std::string &name)
57 meta_kernel_variable(const meta_kernel_variable &other)
58 : m_name(other.m_name)
62 meta_kernel_variable& operator=(const meta_kernel_variable &other)
65 m_name = other.m_name;
71 ~meta_kernel_variable()
75 std::string name() const
85 class meta_kernel_literal
88 typedef T result_type;
90 meta_kernel_literal(const T &value)
95 meta_kernel_literal(const meta_kernel_literal &other)
96 : m_value(other.m_value)
100 meta_kernel_literal& operator=(const meta_kernel_literal &other)
103 m_value = other.m_value;
109 ~meta_kernel_literal()
113 const T& value() const
122 struct meta_kernel_stored_arg
124 meta_kernel_stored_arg()
130 meta_kernel_stored_arg(const meta_kernel_stored_arg &other)
134 set_value(other.m_size, other.m_value);
137 meta_kernel_stored_arg& operator=(const meta_kernel_stored_arg &other)
140 set_value(other.m_size, other.m_value);
147 meta_kernel_stored_arg(const T &value)
154 ~meta_kernel_stored_arg()
161 void set_value(size_t size, const void *value)
170 m_value = std::malloc(size);
171 std::memcpy(m_value, value, size);
179 void set_value(const T &value)
181 set_value(sizeof(T), boost::addressof(value));
188 struct meta_kernel_buffer_info
190 meta_kernel_buffer_info(const buffer &buffer,
191 const std::string &id,
192 memory_object::address_space addr_space,
194 : m_mem(buffer.get()),
196 address_space(addr_space),
202 std::string identifier;
203 memory_object::address_space address_space;
207 struct meta_kernel_svm_info
210 meta_kernel_svm_info(const svm_ptr<T> ptr,
211 const std::string &id,
212 memory_object::address_space addr_space,
216 address_space(addr_space),
223 std::string identifier;
224 memory_object::address_space address_space;
232 struct inject_type_impl
234 void operator()(meta_kernel &)
236 // default implementation does nothing
240 #define BOOST_COMPUTE_META_KERNEL_DECLARE_SCALAR_TYPE_STREAM_OPERATOR(type) \
241 meta_kernel& operator<<(const type &x) \
247 #define BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(type) \
248 meta_kernel& operator<<(const type &x) \
250 m_source << "(" << type_name<type>() << ")"; \
252 for(size_t i = 0; i < vector_size<type>::value; i++){ \
253 *this << lit(x[i]); \
255 if(i != vector_size<type>::value - 1){ \
263 #define BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(type) \
264 BOOST_COMPUTE_META_KERNEL_DECLARE_SCALAR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(type, _)) \
265 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 2), _)) \
266 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 4), _)) \
267 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 8), _)) \
268 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(BOOST_PP_CAT(BOOST_PP_CAT(type, 16), _))
277 argument(const std::string &name, size_t index)
283 const std::string &name() const
298 explicit meta_kernel(const std::string &name)
303 meta_kernel(const meta_kernel &other)
305 m_source.str(other.m_source.str());
306 m_options = other.m_options;
309 meta_kernel& operator=(const meta_kernel &other)
312 m_source.str(other.m_source.str());
313 m_options = other.m_options;
323 std::string name() const
328 std::string source() const
330 std::stringstream stream;
333 if(!m_pragmas.empty()){
334 stream << m_pragmas << "\n";
338 stream << "#define boost_pair_type(t1, t2) _pair_ ## t1 ## _ ## t2 ## _t\n";
339 stream << "#define boost_pair_get(x, n) (n == 0 ? x.first ## x.second)\n";
340 stream << "#define boost_make_pair(t1, x, t2, y) (boost_pair_type(t1, t2)) { x, y }\n";
341 stream << "#define boost_tuple_get(x, n) (x.v ## n)\n";
343 // add type declaration source
344 stream << m_type_declaration_source.str() << "\n";
346 // add external function source
347 stream << m_external_function_source.str() << "\n";
350 stream << "__kernel void " << m_name
351 << "(" << boost::join(m_args, ", ") << ")\n"
352 << "{\n" << m_source.str() << "\n}\n";
357 kernel compile(const context &context, const std::string &options = std::string())
359 // generate the program source
360 std::string source = this->source();
362 // generate cache key
363 std::string cache_key = "__boost_meta_kernel_" +
364 static_cast<std::string>(detail::sha1(source));
366 // load program cache
367 boost::shared_ptr<program_cache> cache =
368 program_cache::get_global_cache(context);
370 std::string compile_options = m_options + options;
372 // load (or build) program from cache
373 ::boost::compute::program program =
374 cache->get_or_build(cache_key, compile_options, source, context);
377 ::boost::compute::kernel kernel = program.create_kernel(name());
380 for(size_t i = 0; i < m_stored_args.size(); i++){
381 const detail::meta_kernel_stored_arg &arg = m_stored_args[i];
384 kernel.set_arg(i, arg.m_size, arg.m_value);
389 for(size_t i = 0; i < m_stored_buffers.size(); i++){
390 const detail::meta_kernel_buffer_info &bi = m_stored_buffers[i];
392 kernel.set_arg(bi.index, bi.m_mem);
396 for(size_t i = 0; i < m_stored_svm_ptrs.size(); i++){
397 const detail::meta_kernel_svm_info &spi = m_stored_svm_ptrs[i];
399 kernel.set_arg_svm_ptr(spi.index, spi.ptr);
406 size_t add_arg(const std::string &name)
408 std::stringstream stream;
409 stream << type<T>() << " " << name;
411 // add argument to list
412 m_args.push_back(stream.str());
415 return m_args.size() - 1;
419 size_t add_arg(memory_object::address_space address_space,
420 const std::string &name)
422 return add_arg_with_qualifiers<T>(address_space_prefix(address_space), name);
426 void set_arg(size_t index, const T &value)
428 if(index >= m_stored_args.size()){
429 m_stored_args.resize(index + 1);
432 m_stored_args[index] = detail::meta_kernel_stored_arg(value);
435 void set_arg(size_t index, const memory_object &mem)
437 set_arg<cl_mem>(index, mem.get());
440 void set_arg(size_t index, const image_sampler &sampler)
442 set_arg<cl_sampler>(index, cl_sampler(sampler));
446 size_t add_set_arg(const std::string &name, const T &value)
448 size_t index = add_arg<T>(name);
449 set_arg<T>(index, value);
453 void add_extension_pragma(const std::string &extension,
454 const std::string &value = "enable")
456 m_pragmas += "#pragma OPENCL EXTENSION " + extension + " : " + value + "\n";
459 void add_extension_pragma(const std::string &extension,
460 const std::string &value) const
462 return const_cast<meta_kernel *>(this)->add_extension_pragma(extension, value);
466 std::string type() const
468 std::stringstream stream;
471 if(boost::is_const<T>::value){
475 // volatile qualifier
476 if(boost::is_volatile<T>::value){
477 stream << "volatile ";
482 typename boost::remove_cv<
483 typename boost::remove_pointer<T>::type
485 stream << type_name<Type>();
488 if(boost::is_pointer<T>::value){
492 // inject type pragmas and/or definitions
499 std::string decl(const std::string &name) const
501 return type<T>() + " " + name;
504 template<class T, class Expr>
505 std::string decl(const std::string &name, const Expr &init) const
507 meta_kernel tmp((std::string()));
508 tmp << tmp.decl<T>(name) << " = " << init;
509 return tmp.m_source.str();
513 detail::meta_kernel_variable<T> var(const std::string &name) const
517 return make_var<T>(name);
521 detail::meta_kernel_literal<T> lit(const T &value) const
525 return detail::meta_kernel_literal<T>(value);
529 detail::meta_kernel_variable<T> expr(const std::string &expr) const
533 return detail::meta_kernel_variable<T>(expr);
536 // define stream operators for scalar and vector types
537 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(char)
538 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(uchar)
539 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(short)
540 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(ushort)
541 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(int)
542 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(uint)
543 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(long)
544 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(ulong)
545 BOOST_COMPUTE_META_KERNEL_DECLARE_TYPE_STREAM_OPERATORS(double)
547 // define stream operators for float scalar and vector types
548 meta_kernel& operator<<(const float &x)
550 m_source << std::showpoint << x << 'f';
554 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float2_)
555 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float4_)
556 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float8_)
557 BOOST_COMPUTE_META_KERNEL_DECLARE_VECTOR_TYPE_STREAM_OPERATOR(float16_)
559 // define stream operators for variable types
561 meta_kernel& operator<<(const meta_kernel_variable<T> &variable)
563 return *this << variable.name();
566 // define stream operators for literal types
568 meta_kernel& operator<<(const meta_kernel_literal<T> &literal)
570 return *this << literal.value();
573 meta_kernel& operator<<(const meta_kernel_literal<bool> &literal)
575 return *this << (literal.value() ? "true" : "false");
578 meta_kernel& operator<<(const meta_kernel_literal<char> &literal)
580 const char c = literal.value();
583 // control characters
585 return *this << "'\\0'";
587 return *this << "'\\a'";
589 return *this << "'\\b'";
591 return *this << "'\\t'";
593 return *this << "'\\n'";
595 return *this << "'\\v'";
597 return *this << "'\\f'";
599 return *this << "'\\r'";
601 // characters which need escaping
606 return *this << "'\\" << c << "'";
608 // all other characters
610 return *this << "'" << c << "'";
614 meta_kernel& operator<<(const meta_kernel_literal<signed char> &literal)
616 return *this << lit<char>(literal.value());
619 meta_kernel& operator<<(const meta_kernel_literal<unsigned char> &literal)
621 return *this << uint_(literal.value());
624 // define stream operators for strings
625 meta_kernel& operator<<(char ch)
631 meta_kernel& operator<<(const char *string)
637 meta_kernel& operator<<(const std::string &string)
644 static detail::meta_kernel_variable<T> make_var(const std::string &name)
646 return detail::meta_kernel_variable<T>(name);
650 static detail::meta_kernel_literal<T> make_lit(const T &value)
652 return detail::meta_kernel_literal<T>(value);
656 static detail::meta_kernel_variable<T> make_expr(const std::string &expr)
658 return detail::meta_kernel_variable<T>(expr);
661 event exec(command_queue &queue)
663 return exec_1d(queue, 0, 1);
666 event exec_1d(command_queue &queue,
667 size_t global_work_offset,
668 size_t global_work_size,
669 const wait_list &events = wait_list())
671 const context &context = queue.get_context();
673 ::boost::compute::kernel kernel = compile(context);
675 return queue.enqueue_1d_range_kernel(
684 event exec_1d(command_queue &queue,
685 size_t global_work_offset,
686 size_t global_work_size,
687 size_t local_work_size,
688 const wait_list &events = wait_list())
690 const context &context = queue.get_context();
692 ::boost::compute::kernel kernel = compile(context);
694 return queue.enqueue_1d_range_kernel(
704 std::string get_buffer_identifier(const buffer &buffer,
705 const memory_object::address_space address_space =
706 memory_object::global_memory)
708 // check if we've already seen buffer
709 for(size_t i = 0; i < m_stored_buffers.size(); i++){
710 const detail::meta_kernel_buffer_info &bi = m_stored_buffers[i];
712 if(bi.m_mem == buffer.get() &&
713 bi.address_space == address_space){
714 return bi.identifier;
718 // create a new binding
719 std::string identifier =
720 "_buf" + lexical_cast<std::string>(m_stored_buffers.size());
721 size_t index = add_arg<T *>(address_space, identifier);
723 // store new buffer info
724 m_stored_buffers.push_back(
725 detail::meta_kernel_buffer_info(buffer, identifier, address_space, index));
731 std::string get_svm_identifier(const svm_ptr<T> &svm_ptr,
732 const memory_object::address_space address_space =
733 memory_object::global_memory)
736 (address_space == memory_object::global_memory)
737 || (address_space == memory_object::constant_memory)
740 // check if we've already seen this pointer
741 for(size_t i = 0; i < m_stored_svm_ptrs.size(); i++){
742 const detail::meta_kernel_svm_info &spi = m_stored_svm_ptrs[i];
744 if(spi.ptr == svm_ptr.get() &&
745 spi.address_space == address_space){
746 return spi.identifier;
750 // create a new binding
751 std::string identifier =
752 "_svm_ptr" + lexical_cast<std::string>(m_stored_svm_ptrs.size());
753 size_t index = add_arg<T *>(address_space, identifier);
755 if(m_stored_svm_ptrs.empty()) {
756 m_options += std::string(" -cl-std=CL2.0");
759 // store new svm pointer info
760 m_stored_svm_ptrs.push_back(
761 detail::meta_kernel_svm_info(
762 svm_ptr, identifier, address_space, index
769 std::string get_image_identifier(const char *qualifiers, const image2d &image)
771 size_t index = add_arg_with_qualifiers<image2d>(qualifiers, "image");
773 set_arg(index, image);
778 std::string get_sampler_identifier(bool normalized_coords,
779 cl_addressing_mode addressing_mode,
780 cl_filter_mode filter_mode)
782 (void) normalized_coords;
783 (void) addressing_mode;
786 m_pragmas += "const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |\n"
787 " CLK_ADDRESS_NONE |\n"
788 " CLK_FILTER_NEAREST;\n";
794 static std::string expr_to_string(const Expr &expr)
796 meta_kernel tmp((std::string()));
798 return tmp.m_source.str();
801 template<class Predicate>
802 detail::invoked_function<bool, boost::tuple<Predicate> > if_(Predicate pred) const
804 return detail::invoked_function<bool, boost::tuple<Predicate> >(
805 "if", std::string(), boost::make_tuple(pred)
809 template<class Predicate>
810 detail::invoked_function<bool, boost::tuple<Predicate> > else_if_(Predicate pred) const
812 return detail::invoked_function<bool, boost::tuple<Predicate> >(
813 "else if", std::string(), boost::make_tuple(pred)
817 detail::meta_kernel_variable<cl_uint> get_global_id(size_t dim) const
819 return expr<cl_uint>("get_global_id(" + lexical_cast<std::string>(dim) + ")");
822 void add_function(const std::string &name, const std::string &source)
824 if(m_external_function_names.count(name)){
828 m_external_function_names.insert(name);
829 m_external_function_source << source << "\n";
832 void add_function(const std::string &name,
833 const std::string &source,
834 const std::map<std::string, std::string> &definitions)
836 typedef std::map<std::string, std::string>::const_iterator iter;
841 for(iter i = definitions.begin(); i != definitions.end(); i++){
842 s << "#define " << i->first;
843 if(!i->second.empty()){
844 s << " " << i->second;
852 for(iter i = definitions.begin(); i != definitions.end(); i++){
853 s << "#undef " << i->first << "\n";
856 add_function(name, s.str());
860 void add_type_declaration(const std::string &declaration)
862 const char *name = type_name<Type>();
864 // check if the type has already been declared
865 std::string source = m_type_declaration_source.str();
866 if(source.find(name) != std::string::npos){
870 m_type_declaration_source << declaration;
874 void inject_type() const
876 inject_type_impl<Type>()(const_cast<meta_kernel &>(*this));
879 // the insert_function_call() method inserts a call to a function with
880 // the given name tuple of argument values.
881 template<class ArgTuple>
882 void insert_function_call(const std::string &name, const ArgTuple &args)
884 *this << name << '(';
885 insert_function_call_args(args);
889 // the insert_function_call_args() method takes a tuple of argument values
890 // and inserts them into the source string with a comma in-between each.
891 // this is useful for creating function calls given a tuple of values.
892 void insert_function_call_args(const boost::tuple<>&)
896 #define BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE(z, n, unused) \
897 inject_type<BOOST_PP_CAT(T, n)>();
899 #define BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG(z, n, unused) \
900 << boost::get<BOOST_PP_DEC(n)>(args) << ", "
902 #define BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS(z, n, unused) \
903 template<BOOST_PP_ENUM_PARAMS(n, class T)> \
904 void insert_function_call_args( \
905 const boost::tuple<BOOST_PP_ENUM_PARAMS(n, T)> &args \
908 BOOST_PP_REPEAT_FROM_TO( \
909 0, n, BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE, ~ \
912 BOOST_PP_REPEAT_FROM_TO( \
913 1, n, BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG, ~ \
915 << boost::get<BOOST_PP_DEC(n)>(args); \
918 BOOST_PP_REPEAT_FROM_TO(
919 1, BOOST_COMPUTE_MAX_ARITY, BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS, ~
922 #undef BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE
923 #undef BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG
924 #undef BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS
926 static const char* address_space_prefix(const memory_object::address_space value)
929 case memory_object::global_memory: return "__global";
930 case memory_object::local_memory: return "__local";
931 case memory_object::private_memory: return "__private";
932 case memory_object::constant_memory: return "__constant";
935 return 0; // unreachable
940 size_t add_arg_with_qualifiers(const char *qualifiers, const std::string &name)
942 size_t index = add_arg<T>(name);
944 // update argument type declaration with qualifiers
946 s << qualifiers << " " << m_args[index];
947 m_args[index] = s.str();
954 std::stringstream m_source;
955 std::stringstream m_external_function_source;
956 std::stringstream m_type_declaration_source;
957 std::set<std::string> m_external_function_names;
958 std::vector<std::string> m_args;
959 std::string m_pragmas;
960 std::string m_options;
961 std::vector<detail::meta_kernel_stored_arg> m_stored_args;
962 std::vector<detail::meta_kernel_buffer_info> m_stored_buffers;
963 std::vector<detail::meta_kernel_svm_info> m_stored_svm_ptrs;
966 template<class ResultType, class ArgTuple>
968 operator<<(meta_kernel &kernel, const invoked_function<ResultType, ArgTuple> &expr)
970 if(!expr.source().empty()){
971 kernel.add_function(expr.name(), expr.source(), expr.definitions());
974 kernel.insert_function_call(expr.name(), expr.args());
979 template<class ResultType, class ArgTuple, class CaptureTuple>
981 operator<<(meta_kernel &kernel,
982 const invoked_closure<ResultType, ArgTuple, CaptureTuple> &expr)
984 if(!expr.source().empty()){
985 kernel.add_function(expr.name(), expr.source(), expr.definitions());
988 kernel << expr.name() << '(';
989 kernel.insert_function_call_args(expr.args());
991 kernel.insert_function_call_args(expr.capture());
997 template<class Arg1, class Arg2, class Result>
998 inline meta_kernel& operator<<(meta_kernel &kernel,
999 const invoked_binary_operator<Arg1,
1003 return kernel << "((" << expr.arg1() << ")"
1005 << "(" << expr.arg2() << "))";
1008 template<class T, class IndexExpr>
1009 inline meta_kernel& operator<<(meta_kernel &kernel,
1010 const detail::device_ptr_index_expr<T, IndexExpr> &expr)
1012 if(expr.m_index == 0){
1014 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1015 '[' << expr.m_expr << ']';
1019 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1020 '[' << expr.m_index << "+(" << expr.m_expr << ")]";
1024 template<class T1, class T2, class IndexExpr>
1025 inline meta_kernel& operator<<(meta_kernel &kernel,
1026 const detail::device_ptr_index_expr<std::pair<T1, T2>, IndexExpr> &expr)
1028 typedef std::pair<T1, T2> T;
1030 if(expr.m_index == 0){
1032 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1033 '[' << expr.m_expr << ']';
1037 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1038 '[' << expr.m_index << "+(" << expr.m_expr << ")]";
1042 // SVM requires OpenCL 2.0
1043 #if defined(BOOST_COMPUTE_CL_VERSION_2_0) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1044 template<class T, class IndexExpr>
1045 inline meta_kernel& operator<<(meta_kernel &kernel,
1046 const svm_ptr_index_expr<T, IndexExpr> &expr)
1049 kernel.get_svm_identifier<T>(expr.m_svm_ptr) <<
1050 '[' << expr.m_expr << ']';
1054 template<class Predicate, class Arg>
1055 inline meta_kernel& operator<<(meta_kernel &kernel,
1056 const invoked_unary_negate_function<Predicate,
1059 return kernel << "!(" << expr.pred()(expr.expr()) << ')';
1062 template<class Predicate, class Arg1, class Arg2>
1063 inline meta_kernel& operator<<(meta_kernel &kernel,
1064 const invoked_binary_negate_function<Predicate,
1068 return kernel << "!(" << expr.pred()(expr.expr1(), expr.expr2()) << ')';
1071 // get<N>() for vector types
1072 template<size_t N, class Arg, class T>
1073 inline meta_kernel& operator<<(meta_kernel &kernel,
1074 const invoked_get<N, Arg, T> &expr)
1076 BOOST_STATIC_ASSERT(N < 16);
1079 return kernel << expr.m_arg << ".s" << int_(N);
1083 # pragma warning(push)
1084 # pragma warning(disable: 4307)
1086 return kernel << expr.m_arg << ".s" << char('a' + (N - 10));
1088 # pragma warning(pop)
1095 template<class T, class Arg>
1096 inline meta_kernel& operator<<(meta_kernel &kernel,
1097 const invoked_field<T, Arg> &expr)
1099 return kernel << expr.m_arg << "." << expr.m_field;
1102 template<class T, class Arg>
1103 inline meta_kernel& operator<<(meta_kernel &k,
1104 const invoked_as<T, Arg> &expr)
1106 return k << "as_" << type_name<T>() << "(" << expr.m_arg << ")";
1109 template<class T, class Arg>
1110 inline meta_kernel& operator<<(meta_kernel &k,
1111 const invoked_convert<T, Arg> &expr)
1113 return k << "convert_" << type_name<T>() << "(" << expr.m_arg << ")";
1116 template<class T, class Arg>
1117 inline meta_kernel& operator<<(meta_kernel &k,
1118 const invoked_identity<T, Arg> &expr)
1120 return k << expr.m_arg;
1124 struct inject_type_impl<double_>
1126 void operator()(meta_kernel &kernel)
1128 kernel.add_extension_pragma("cl_khr_fp64", "enable");
1132 template<class Scalar, size_t N>
1133 struct inject_type_impl<vector_type<Scalar, N> >
1135 void operator()(meta_kernel &kernel)
1137 kernel.inject_type<Scalar>();
1141 } // end detail namespace
1142 } // end compute namespace
1143 } // end boost namespace
1145 #endif // BOOST_COMPUTE_DETAIL_META_KERNEL_HPP