Name	Size	Mode	Actions
balanced_path.hpp	6267	0644	editdlrm
binary_find.hpp	4626	0644	editdlrm
compact.hpp	2188	0644	editdlrm
copy_on_device.hpp	7382	0644	editdlrm
copy_to_device.hpp	6731	0644	editdlrm
copy_to_host.hpp	6658	0644	editdlrm
count_if_with_ballot.hpp	2605	0644	editdlrm
count_if_with_reduce.hpp	2648	0644	editdlrm
count_if_with_threads.hpp	4241	0644	editdlrm
find_extrema.hpp	2441	0644	editdlrm
find_extrema_on_cpu.hpp	5333	0644	editdlrm
find_extrema_with_atomics.hpp	4251	0644	editdlrm
find_extrema_with_reduce.hpp	18672	0644	editdlrm
find_if_with_atomics.hpp	8509	0644	editdlrm
inplace_reduce.hpp	4839	0644	editdlrm
insertion_sort.hpp	6030	0644	editdlrm
merge_path.hpp	3874	0644	editdlrm
merge_sort_on_cpu.hpp	14654	0644	editdlrm
merge_sort_on_gpu.hpp	22093	0644	editdlrm
merge_with_merge_path.hpp	7339	0644	editdlrm
radix_sort.hpp	15502	0644	editdlrm
random_fill.hpp	1880	0644	editdlrm
reduce_by_key.hpp	5243	0644	editdlrm
reduce_by_key_with_scan.hpp	23735	0644	editdlrm
reduce_on_cpu.hpp	3949	0644	editdlrm
reduce_on_gpu.hpp	10391	0644	editdlrm
scan.hpp	1545	0644	editdlrm
scan_on_cpu.hpp	7056	0644	editdlrm
scan_on_gpu.hpp	11320	0644	editdlrm
search_all.hpp	2579	0644	editdlrm
serial_accumulate.hpp	1990	0644	editdlrm
serial_count_if.hpp	2243	0644	editdlrm
serial_find_extrema.hpp	3151	0644	editdlrm
serial_merge.hpp	3504	0644	editdlrm
serial_reduce.hpp	2145	0644	editdlrm
serial_reduce_by_key.hpp	4240	0644	editdlrm
serial_scan.hpp	3184	0644	editdlrm

//---------------------------------------------------------------------------// // Copyright (c) 2013 Kyle Lutz // // Distributed under the Boost Software License, Version 1.0 // See accompanying file LICENSE_1_0.txt or copy at // http://www.boost.org/LICENSE_1_0.txt // // See http://boostorg.github.com/compute for more information. //---------------------------------------------------------------------------// #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP #define BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP #include #include #include #include #include #include namespace boost { namespace compute { namespace detail { template inline InputIterator find_extrema_with_atomics(InputIterator first, InputIterator last, Compare compare, const bool find_minimum, command_queue &queue) { typedef typename std::iterator_traits::value_type value_type; typedef typename std::iterator_traits::difference_type difference_type; const context &context = queue.get_context(); meta_kernel k("find_extrema"); atomic_cmpxchg atomic_cmpxchg_uint; k << "const uint gid = get_global_id(0);\n" << "uint old_index = *index;\n" << k.decl("old") << " = " << first[k.var("old_index")] << ";\n" << k.decl("new") << " = " << first[k.var("gid")] << ";\n" << k.decl("compare_result") << ";\n" << "#ifdef BOOST_COMPUTE_FIND_MAXIMUM\n" << "while(" << "(compare_result = " << compare(k.var("old"), k.var("new")) << ")" << " || (!(compare_result" << " || " << compare(k.var("new"), k.var("old")) << ") " "&& gid < old_index)){\n" << "#else\n" << // while condition explained for minimum case with less (<) // as comparison function: // while(new_value < old_value // OR (new_value == old_value AND new_index < old_index)) "while(" << "(compare_result = " << compare(k.var("new"), k.var("old")) << ")" << " || (!(compare_result" << " || " << compare(k.var("old"), k.var("new")) << ") " "&& gid < old_index)){\n" << "#endif\n" << " if(" << atomic_cmpxchg_uint(k.var("index"), k.var("old_index"), k.var("gid")) << " == old_index)\n" << " break;\n" << " else\n" << " old_index = *index;\n" << "old = " << first[k.var("old_index")] << ";\n" << "}\n"; size_t index_arg_index = k.add_arg(memory_object::global_memory, "index"); std::string options; if(!find_minimum){ options = "-DBOOST_COMPUTE_FIND_MAXIMUM"; } kernel kernel = k.compile(context, options); // setup index buffer scalar index(context); kernel.set_arg(index_arg_index, index.get_buffer()); // initialize index index.write(0, queue); // run kernel size_t count = iterator_range_size(first, last); queue.enqueue_1d_range_kernel(kernel, 0, count, 0); // read index and return iterator return first + static_cast(index.read(queue)); } } // end detail namespace } // end compute namespace } // end boost namespace #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_FIND_EXTREMA_WITH_ATOMICS_HPP