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)
670 const context &context = queue.get_context();
672 ::boost::compute::kernel kernel = compile(context);
674 return queue.enqueue_1d_range_kernel(
682 event exec_1d(command_queue &queue,
683 size_t global_work_offset,
684 size_t global_work_size,
685 size_t local_work_size)
687 const context &context = queue.get_context();
689 ::boost::compute::kernel kernel = compile(context);
691 return queue.enqueue_1d_range_kernel(
700 std::string get_buffer_identifier(const buffer &buffer,
701 const memory_object::address_space address_space =
702 memory_object::global_memory)
704 // check if we've already seen buffer
705 for(size_t i = 0; i < m_stored_buffers.size(); i++){
706 const detail::meta_kernel_buffer_info &bi = m_stored_buffers[i];
708 if(bi.m_mem == buffer.get() &&
709 bi.address_space == address_space){
710 return bi.identifier;
714 // create a new binding
715 std::string identifier =
716 "_buf" + lexical_cast<std::string>(m_stored_buffers.size());
717 size_t index = add_arg<T *>(address_space, identifier);
719 // store new buffer info
720 m_stored_buffers.push_back(
721 detail::meta_kernel_buffer_info(buffer, identifier, address_space, index));
727 std::string get_svm_identifier(const svm_ptr<T> &svm_ptr,
728 const memory_object::address_space address_space =
729 memory_object::global_memory)
732 (address_space == memory_object::global_memory)
733 || (address_space == memory_object::constant_memory)
736 // check if we've already seen this pointer
737 for(size_t i = 0; i < m_stored_svm_ptrs.size(); i++){
738 const detail::meta_kernel_svm_info &spi = m_stored_svm_ptrs[i];
740 if(spi.ptr == svm_ptr.get() &&
741 spi.address_space == address_space){
742 return spi.identifier;
746 // create a new binding
747 std::string identifier =
748 "_svm_ptr" + lexical_cast<std::string>(m_stored_svm_ptrs.size());
749 size_t index = add_arg<T *>(address_space, identifier);
751 if(m_stored_svm_ptrs.empty()) {
752 m_options += std::string(" -cl-std=CL2.0");
755 // store new svm pointer info
756 m_stored_svm_ptrs.push_back(
757 detail::meta_kernel_svm_info(
758 svm_ptr, identifier, address_space, index
765 std::string get_image_identifier(const char *qualifiers, const image2d &image)
767 size_t index = add_arg_with_qualifiers<image2d>(qualifiers, "image");
769 set_arg(index, image);
774 std::string get_sampler_identifier(bool normalized_coords,
775 cl_addressing_mode addressing_mode,
776 cl_filter_mode filter_mode)
778 (void) normalized_coords;
779 (void) addressing_mode;
782 m_pragmas += "const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE |\n"
783 " CLK_ADDRESS_NONE |\n"
784 " CLK_FILTER_NEAREST;\n";
790 static std::string expr_to_string(const Expr &expr)
792 meta_kernel tmp((std::string()));
794 return tmp.m_source.str();
797 template<class Predicate>
798 detail::invoked_function<bool, boost::tuple<Predicate> > if_(Predicate pred) const
800 return detail::invoked_function<bool, boost::tuple<Predicate> >(
801 "if", std::string(), boost::make_tuple(pred)
805 template<class Predicate>
806 detail::invoked_function<bool, boost::tuple<Predicate> > else_if_(Predicate pred) const
808 return detail::invoked_function<bool, boost::tuple<Predicate> >(
809 "else if", std::string(), boost::make_tuple(pred)
813 detail::meta_kernel_variable<cl_uint> get_global_id(size_t dim) const
815 return expr<cl_uint>("get_global_id(" + lexical_cast<std::string>(dim) + ")");
818 void add_function(const std::string &name, const std::string &source)
820 if(m_external_function_names.count(name)){
824 m_external_function_names.insert(name);
825 m_external_function_source << source << "\n";
828 void add_function(const std::string &name,
829 const std::string &source,
830 const std::map<std::string, std::string> &definitions)
832 typedef std::map<std::string, std::string>::const_iterator iter;
837 for(iter i = definitions.begin(); i != definitions.end(); i++){
838 s << "#define " << i->first;
839 if(!i->second.empty()){
840 s << " " << i->second;
848 for(iter i = definitions.begin(); i != definitions.end(); i++){
849 s << "#undef " << i->first << "\n";
852 add_function(name, s.str());
856 void add_type_declaration(const std::string &declaration)
858 const char *name = type_name<Type>();
860 // check if the type has already been declared
861 std::string source = m_type_declaration_source.str();
862 if(source.find(name) != std::string::npos){
866 m_type_declaration_source << declaration;
870 void inject_type() const
872 inject_type_impl<Type>()(const_cast<meta_kernel &>(*this));
875 // the insert_function_call() method inserts a call to a function with
876 // the given name tuple of argument values.
877 template<class ArgTuple>
878 void insert_function_call(const std::string &name, const ArgTuple &args)
880 *this << name << '(';
881 insert_function_call_args(args);
885 // the insert_function_call_args() method takes a tuple of argument values
886 // and inserts them into the source string with a comma in-between each.
887 // this is useful for creating function calls given a tuple of values.
888 void insert_function_call_args(const boost::tuple<>&)
892 #define BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE(z, n, unused) \
893 inject_type<BOOST_PP_CAT(T, n)>();
895 #define BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG(z, n, unused) \
896 << boost::get<BOOST_PP_DEC(n)>(args) << ", "
898 #define BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS(z, n, unused) \
899 template<BOOST_PP_ENUM_PARAMS(n, class T)> \
900 void insert_function_call_args( \
901 const boost::tuple<BOOST_PP_ENUM_PARAMS(n, T)> &args \
904 BOOST_PP_REPEAT_FROM_TO( \
905 0, n, BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE, ~ \
908 BOOST_PP_REPEAT_FROM_TO( \
909 1, n, BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG, ~ \
911 << boost::get<BOOST_PP_DEC(n)>(args); \
914 BOOST_PP_REPEAT_FROM_TO(
915 1, BOOST_COMPUTE_MAX_ARITY, BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS, ~
918 #undef BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARG_TYPE
919 #undef BOOST_COMPUTE_META_KERNEL_STREAM_FUNCTION_ARG
920 #undef BOOST_COMPUTE_META_KERNEL_INSERT_FUNCTION_ARGS
922 static const char* address_space_prefix(const memory_object::address_space value)
925 case memory_object::global_memory: return "__global";
926 case memory_object::local_memory: return "__local";
927 case memory_object::private_memory: return "__private";
928 case memory_object::constant_memory: return "__constant";
931 return 0; // unreachable
936 size_t add_arg_with_qualifiers(const char *qualifiers, const std::string &name)
938 size_t index = add_arg<T>(name);
940 // update argument type declaration with qualifiers
942 s << qualifiers << " " << m_args[index];
943 m_args[index] = s.str();
950 std::stringstream m_source;
951 std::stringstream m_external_function_source;
952 std::stringstream m_type_declaration_source;
953 std::set<std::string> m_external_function_names;
954 std::vector<std::string> m_args;
955 std::string m_pragmas;
956 std::string m_options;
957 std::vector<detail::meta_kernel_stored_arg> m_stored_args;
958 std::vector<detail::meta_kernel_buffer_info> m_stored_buffers;
959 std::vector<detail::meta_kernel_svm_info> m_stored_svm_ptrs;
962 template<class ResultType, class ArgTuple>
964 operator<<(meta_kernel &kernel, const invoked_function<ResultType, ArgTuple> &expr)
966 if(!expr.source().empty()){
967 kernel.add_function(expr.name(), expr.source(), expr.definitions());
970 kernel.insert_function_call(expr.name(), expr.args());
975 template<class ResultType, class ArgTuple, class CaptureTuple>
977 operator<<(meta_kernel &kernel,
978 const invoked_closure<ResultType, ArgTuple, CaptureTuple> &expr)
980 if(!expr.source().empty()){
981 kernel.add_function(expr.name(), expr.source(), expr.definitions());
984 kernel << expr.name() << '(';
985 kernel.insert_function_call_args(expr.args());
987 kernel.insert_function_call_args(expr.capture());
993 template<class Arg1, class Arg2, class Result>
994 inline meta_kernel& operator<<(meta_kernel &kernel,
995 const invoked_binary_operator<Arg1,
999 return kernel << "((" << expr.arg1() << ")"
1001 << "(" << expr.arg2() << "))";
1004 template<class T, class IndexExpr>
1005 inline meta_kernel& operator<<(meta_kernel &kernel,
1006 const detail::device_ptr_index_expr<T, IndexExpr> &expr)
1008 if(expr.m_index == 0){
1010 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1011 '[' << expr.m_expr << ']';
1015 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1016 '[' << expr.m_index << "+(" << expr.m_expr << ")]";
1020 template<class T1, class T2, class IndexExpr>
1021 inline meta_kernel& operator<<(meta_kernel &kernel,
1022 const detail::device_ptr_index_expr<std::pair<T1, T2>, IndexExpr> &expr)
1024 typedef std::pair<T1, T2> T;
1026 if(expr.m_index == 0){
1028 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1029 '[' << expr.m_expr << ']';
1033 kernel.get_buffer_identifier<T>(expr.m_buffer) <<
1034 '[' << expr.m_index << "+(" << expr.m_expr << ")]";
1038 // SVM requires OpenCL 2.0
1039 #if defined(CL_VERSION_2_0) || defined(BOOST_COMPUTE_DOXYGEN_INVOKED)
1040 template<class T, class IndexExpr>
1041 inline meta_kernel& operator<<(meta_kernel &kernel,
1042 const svm_ptr_index_expr<T, IndexExpr> &expr)
1045 kernel.get_svm_identifier<T>(expr.m_svm_ptr) <<
1046 '[' << expr.m_expr << ']';
1050 template<class Predicate, class Arg>
1051 inline meta_kernel& operator<<(meta_kernel &kernel,
1052 const invoked_unary_negate_function<Predicate,
1055 return kernel << "!(" << expr.pred()(expr.expr()) << ')';
1058 template<class Predicate, class Arg1, class Arg2>
1059 inline meta_kernel& operator<<(meta_kernel &kernel,
1060 const invoked_binary_negate_function<Predicate,
1064 return kernel << "!(" << expr.pred()(expr.expr1(), expr.expr2()) << ')';
1067 // get<N>() for vector types
1068 template<size_t N, class Arg, class T>
1069 inline meta_kernel& operator<<(meta_kernel &kernel,
1070 const invoked_get<N, Arg, T> &expr)
1072 BOOST_STATIC_ASSERT(N < 16);
1075 return kernel << expr.m_arg << ".s" << uint_(N);
1079 # pragma warning(push)
1080 # pragma warning(disable: 4307)
1082 return kernel << expr.m_arg << ".s" << char('a' + (N - 10));
1084 # pragma warning(pop)
1091 template<class T, class Arg>
1092 inline meta_kernel& operator<<(meta_kernel &kernel,
1093 const invoked_field<T, Arg> &expr)
1095 return kernel << expr.m_arg << "." << expr.m_field;
1098 template<class T, class Arg>
1099 inline meta_kernel& operator<<(meta_kernel &k,
1100 const invoked_as<T, Arg> &expr)
1102 return k << "as_" << type_name<T>() << "(" << expr.m_arg << ")";
1105 template<class T, class Arg>
1106 inline meta_kernel& operator<<(meta_kernel &k,
1107 const invoked_convert<T, Arg> &expr)
1109 return k << "convert_" << type_name<T>() << "(" << expr.m_arg << ")";
1112 template<class T, class Arg>
1113 inline meta_kernel& operator<<(meta_kernel &k,
1114 const invoked_identity<T, Arg> &expr)
1116 return k << expr.m_arg;
1120 struct inject_type_impl<double_>
1122 void operator()(meta_kernel &kernel)
1124 kernel.add_extension_pragma("cl_khr_fp64", "enable");
1128 template<class Scalar, size_t N>
1129 struct inject_type_impl<vector_type<Scalar, N> >
1131 void operator()(meta_kernel &kernel)
1133 kernel.inject_type<Scalar>();
1137 } // end detail namespace
1138 } // end compute namespace
1139 } // end boost namespace
1141 #endif // BOOST_COMPUTE_DETAIL_META_KERNEL_HPP