projects / ceph.git / blame
| Commit | Line | Data |
|---|---|---|
| 7c673cae FG |
1 | //---------------------------------------------------------------------------// |
| 2 | // Copyright (c) 2014 Roshan <thisisroshansmail@gmail.com> | |
| 3 | // | |
| 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 | |
| 7 | // | |
| 8 | // See http://boostorg.github.com/compute for more information. | |
| 9 | //---------------------------------------------------------------------------// | |
| 10 | ||
| 11 | #ifndef BOOST_COMPUTE_ALGORITHM_DETAIL_COMPACT_HPP | |
| 12 | #define BOOST_COMPUTE_ALGORITHM_DETAIL_COMPACT_HPP | |
| 13 | ||
| 14 | #include <iterator> | |
| 15 | ||
| 16 | #include <boost/compute/container/vector.hpp> | |
| 17 | #include <boost/compute/detail/iterator_range_size.hpp> | |
| 18 | #include <boost/compute/detail/meta_kernel.hpp> | |
| 19 | #include <boost/compute/system.hpp> | |
| 20 | ||
| 21 | namespace boost { | |
| 22 | namespace compute { | |
| 23 | namespace detail { | |
| 24 | ||
| 25 | /// | |
| 26 | /// \brief Compact kernel class | |
| 27 | /// | |
| 28 | /// Subclass of meta_kernel to compact the result of set kernels to | |
| 29 | /// get actual sets | |
| 30 | /// | |
| 31 | class compact_kernel : public meta_kernel | |
| 32 | { | |
| 33 | public: | |
| 34 | unsigned int tile_size; | |
| 35 | ||
| 36 | compact_kernel() : meta_kernel("compact") | |
| 37 | { | |
| 38 | tile_size = 4; | |
| 39 | } | |
| 40 | ||
| 41 | template<class InputIterator1, class InputIterator2, class OutputIterator> | |
| 42 | void set_range(InputIterator1 start, | |
| 43 | InputIterator2 counts_begin, | |
| 44 | InputIterator2 counts_end, | |
| 45 | OutputIterator result) | |
| 46 | { | |
| 47 | m_count = iterator_range_size(counts_begin, counts_end) - 1; | |
| 48 | ||
| 49 | *this << | |
| 50 | "uint i = get_global_id(0);\n" << | |
| 51 | "uint count = i*" << tile_size << ";\n" << | |
| 52 | "for(uint j = " << counts_begin[expr<uint_>("i")] << "; j<" << | |
| 53 | counts_begin[expr<uint_>("i+1")] << "; j++, count++)\n" << | |
| 54 | "{\n" << | |
| 55 | result[expr<uint_>("j")] << " = " << start[expr<uint_>("count")] | |
| 56 | << ";\n" << | |
| 57 | "}\n"; | |
| 58 | } | |
| 59 | ||
| 60 | event exec(command_queue &queue) | |
| 61 | { | |
| 62 | if(m_count == 0) { | |
| 63 | return event(); | |
| 64 | } | |
| 65 | ||
| 66 | return exec_1d(queue, 0, m_count); | |
| 67 | } | |
| 68 | ||
| 69 | private: | |
| 70 | size_t m_count; | |
| 71 | }; | |
| 72 | ||
| 73 | } //end detail namespace | |
| 74 | } //end compute namespace | |
| 75 | } //end boost namespace | |
| 76 | ||
| 77 | #endif // BOOST_COMPUTE_ALGORITHM_DETAIL_COMPACT_HPP |