| 190 | |
| 191 | template <class ForwardIt1, class ForwardIt2, class UnaryPredicate> |
| 192 | sycl::event copy_if(sycl::queue &q, util::allocation_group &scratch_allocations, |
| 193 | ForwardIt1 first, ForwardIt1 last, ForwardIt2 d_first, |
| 194 | UnaryPredicate pred, |
| 195 | std::size_t *num_elements_copied = nullptr, |
| 196 | const std::vector<sycl::event> &deps = {}) { |
| 197 | if(first == last) { |
| 198 | if(num_elements_copied) |
| 199 | *num_elements_copied = 0; |
| 200 | return sycl::event{}; |
| 201 | } |
| 202 | |
| 203 | // TODO: We could optimize by switching between 32/64 bit types |
| 204 | // depending on problem size |
| 205 | using ScanT = std::size_t; |
| 206 | |
| 207 | auto generator = [=](auto idx, auto effective_group_id, |
| 208 | auto effective_global_id, auto problem_size) { |
| 209 | if(effective_global_id >= problem_size) |
| 210 | return ScanT{0}; |
| 211 | |
| 212 | ForwardIt1 it = first; |
| 213 | std::advance(it, effective_global_id); |
| 214 | if(pred(*it)) |
| 215 | return ScanT{1}; |
| 216 | |
| 217 | return ScanT{0}; |
| 218 | }; |
| 219 | |
| 220 | auto result_processor = [=](auto idx, auto effective_group_id, |
| 221 | auto effective_global_id, auto problem_size, |
| 222 | auto value) { |
| 223 | if (effective_global_id < problem_size) { |
| 224 | ForwardIt2 output = d_first; |
| 225 | ForwardIt1 input = first; |
| 226 | std::advance(input, effective_global_id); |
| 227 | std::advance(output, value); |
| 228 | |
| 229 | bool needs_copy = false; |
| 230 | |
| 231 | if(effective_global_id < problem_size) { |
| 232 | auto input_value = *input; |
| 233 | needs_copy = pred(input_value); |
| 234 | if(needs_copy) |
| 235 | *output = *input; |
| 236 | } |
| 237 | |
| 238 | if (effective_global_id == problem_size - 1 && num_elements_copied) { |
| 239 | ScanT inclusive_scan_result = value; |
| 240 | // We did an exclusive scan, so if the last element also was copied, |
| 241 | // we need to add that. |
| 242 | if(needs_copy) |
| 243 | ++inclusive_scan_result; |
| 244 | |
| 245 | *num_elements_copied = static_cast<std::size_t>(inclusive_scan_result); |
| 246 | } |
| 247 | } |
| 248 | }; |
| 249 |
no outgoing calls