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