| 36 | |
| 37 | template<class InputIterator, class OutputIterator, class BinaryFunction> |
| 38 | size_t reduce(InputIterator first, |
| 39 | size_t count, |
| 40 | OutputIterator result, |
| 41 | size_t block_size, |
| 42 | BinaryFunction function, |
| 43 | command_queue &queue) |
| 44 | { |
| 45 | typedef typename |
| 46 | std::iterator_traits<InputIterator>::value_type |
| 47 | input_type; |
| 48 | typedef typename |
| 49 | boost::compute::result_of<BinaryFunction(input_type, input_type)>::type |
| 50 | result_type; |
| 51 | |
| 52 | const context &context = queue.get_context(); |
| 53 | size_t block_count = count / 2 / block_size; |
| 54 | size_t total_block_count = |
| 55 | static_cast<size_t>(std::ceil(float(count) / 2.f / float(block_size))); |
| 56 | |
| 57 | if(block_count != 0){ |
| 58 | meta_kernel k("block_reduce"); |
| 59 | size_t output_arg = k.add_arg<result_type *>(memory_object::global_memory, "output"); |
| 60 | size_t block_arg = k.add_arg<input_type *>(memory_object::local_memory, "block"); |
| 61 | |
| 62 | k << |
| 63 | "const uint gid = get_global_id(0);\n" << |
| 64 | "const uint lid = get_local_id(0);\n" << |
| 65 | |
| 66 | // copy values to local memory |
| 67 | "block[lid] = " << |
| 68 | function(first[k.make_var<uint_>("gid*2+0")], |
| 69 | first[k.make_var<uint_>("gid*2+1")]) << ";\n" << |
| 70 | |
| 71 | // perform reduction |
| 72 | "for(uint i = 1; i < " << uint_(block_size) << "; i <<= 1){\n" << |
| 73 | " barrier(CLK_LOCAL_MEM_FENCE);\n" << |
| 74 | " uint mask = (i << 1) - 1;\n" << |
| 75 | " if((lid & mask) == 0){\n" << |
| 76 | " block[lid] = " << |
| 77 | function(k.expr<input_type>("block[lid]"), |
| 78 | k.expr<input_type>("block[lid+i]")) << ";\n" << |
| 79 | " }\n" << |
| 80 | "}\n" << |
| 81 | |
| 82 | // write block result to global output |
| 83 | "if(lid == 0)\n" << |
| 84 | " output[get_group_id(0)] = block[0];\n"; |
| 85 | |
| 86 | kernel kernel = k.compile(context); |
| 87 | kernel.set_arg(output_arg, result.get_buffer()); |
| 88 | kernel.set_arg(block_arg, local_buffer<input_type>(block_size)); |
| 89 | |
| 90 | queue.enqueue_1d_range_kernel(kernel, |
| 91 | 0, |
| 92 | block_count * block_size, |
| 93 | block_size); |
| 94 | } |
| 95 | |