| 77 | /// |
| 78 | template<class InputIterator, class ValueType> |
| 79 | inline InputIterator pp_floor(InputIterator first, |
| 80 | InputIterator last, |
| 81 | ValueType value, |
| 82 | command_queue &queue) |
| 83 | { |
| 84 | typedef typename std::iterator_traits<InputIterator>::value_type value_type; |
| 85 | |
| 86 | size_t count = detail::iterator_range_size(first, last); |
| 87 | if(count == 0){ |
| 88 | return last; |
| 89 | } |
| 90 | const context &context = queue.get_context(); |
| 91 | |
| 92 | detail::meta_kernel k("pp_floor"); |
| 93 | size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index"); |
| 94 | size_t value_arg = k.add_arg<value_type>(memory_object::private_memory, "value"); |
| 95 | atomic_max<int_> atomic_max_int; |
| 96 | |
| 97 | k << k.decl<const int_>("i") << " = get_global_id(0);\n" |
| 98 | << k.decl<const value_type>("cur_value") << "=" |
| 99 | << first[k.var<const int_>("i")] << ";\n" |
| 100 | << "if(cur_value >= " << first[k.expr<int_>("*index")] |
| 101 | << " && cur_value < value){\n" |
| 102 | << " " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n" |
| 103 | << "}\n"; |
| 104 | |
| 105 | kernel kernel = k.compile(context); |
| 106 | |
| 107 | scalar<int_> index(context); |
| 108 | kernel.set_arg(index_arg, index.get_buffer()); |
| 109 | |
| 110 | index.write(static_cast<int_>(0), queue); |
| 111 | |
| 112 | kernel.set_arg(value_arg, value); |
| 113 | |
| 114 | queue.enqueue_1d_range_kernel(kernel, 0, count, 0); |
| 115 | |
| 116 | int result = static_cast<int>(index.read(queue)); |
| 117 | return first + result; |
| 118 | } |
| 119 | |
| 120 | } // end detail namespace |
| 121 |
no test coverage detected