inline InputIterator pp_floor(InputIterator first,
InputIterator last,
ValueType value,
command_queue &queue)
{
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
size_t count = detail::iterator_range_size(first, last);
if(count == 0){
return last;
}
const context &context = queue.get_context();
detail::meta_kernel k("pp_floor");
size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index");
size_t value_arg = k.add_arg<value_type>(memory_object::private_memory, "value");
atomic_max<int_> atomic_max_int;
k << k.decl<const int_>("i") << " = get_global_id(0);\n"
<< k.decl<const value_type>("cur_value") << "="
<< first[k.var<const int_>("i")] << ";\n"
<< "if(cur_value >= " << first[k.expr<int_>("*index")]
<< " && cur_value < value){\n"
<< " " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
<< "}\n";
kernel kernel = k.compile(context);
scalar<int_> index(context);
kernel.set_arg(index_arg, index.get_buffer());
index.write(static_cast<int_>(0), queue);
kernel.set_arg(value_arg, value);
queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
int result = static_cast<int>(index.read(queue));
return first + result;
}
inline InputIterator prev_permutation_helper(InputIterator first,
InputIterator last,
command_queue &queue)
{
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
size_t count = detail::iterator_range_size(first, last);
if(count == 0 || count == 1){
return last;
}
count = count - 1;
const context &context = queue.get_context();
detail::meta_kernel k("prev_permutation");
size_t index_arg = k.add_arg<int *>(memory_object::global_memory, "index");
atomic_max<int_> atomic_max_int;
k << k.decl<const int_>("i") << " = get_global_id(0);\n"
<< k.decl<const value_type>("cur_value") << "="
<< first[k.var<const int_>("i")] << ";\n"
<< k.decl<const value_type>("next_value") << "="
<< first[k.expr<const int_>("i+1")] << ";\n"
<< "if(cur_value > next_value){\n"
<< " " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
<< "}\n";
kernel kernel = k.compile(context);
scalar<int_> index(context);
kernel.set_arg(index_arg, index.get_buffer());
index.write(static_cast<int_>(-1), queue);
queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
int result = static_cast<int>(index.read(queue));
if(result == -1) return last;
else return first + result;
}
inline InputIterator find_end_helper(InputIterator first,
InputIterator last,
UnaryPredicate predicate,
command_queue &queue)
{
typedef typename std::iterator_traits<InputIterator>::value_type value_type;
size_t count = detail::iterator_range_size(first, last);
if(count == 0){
return last;
}
const context &context = queue.get_context();
detail::meta_kernel k("find_end");
size_t index_arg = k.add_arg<int *>("__global", "index");
atomic_max<int_> atomic_max_int;
k << k.decl<const int_>("i") << " = get_global_id(0);\n"
<< k.decl<const value_type>("value") << "="
<< first[k.var<const int_>("i")] << ";\n"
<< "if(" << predicate(k.var<const value_type>("value")) << "){\n"
<< " " << atomic_max_int(k.var<int_ *>("index"), k.var<int_>("i")) << ";\n"
<< "}\n";
kernel kernel = k.compile(context);
scalar<int_> index(context);
kernel.set_arg(index_arg, index.get_buffer());
index.write(static_cast<int_>(-1), queue);
queue.enqueue_1d_range_kernel(kernel, 0, count, 0);
int result = static_cast<int>(index.read(queue));
if(result == -1) return last;
else return first + result;
}
